fix: use cudaMemcpyDefault for D2D copy (handles CPU->GPU too); force only nodes; add phase2 debug fprintfs

Author: Alex7MV

src/llama-context.cpp

@@ -60,6 +60,7 @@ extern "C" {
     int cudaStreamSynchronize(void * stream);
 }
 constexpr int cudaMemcpyDeviceToHost = 2;
+constexpr int cudaMemcpyDefault = 4;
 constexpr int cudaMemcpyDeviceToDevice = 3;
 
 // Static memory hijacking: additional CUDA API forward declarations
@@ -356,7 +357,7 @@ void phase2_hijack::copy_data_to_static(void* cuda_stream) {
     for (int i = 0; i < n_slots; i++) {
         if (slots[i].addr && slots[i].orig_data && slots[i].size > 0) {
             cudaMemcpyAsync(slots[i].addr, slots[i].orig_data,
-                slots[i].size, cudaMemcpyDeviceToDevice, (cudaStream_t)cuda_stream);
+                slots[i].size, cudaMemcpyDefault, (cudaStream_t)cuda_stream);
         }
     }
 }
@@ -2084,21 +2085,6 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll
                                 ggml_tensor * t = ggml_graph_node(phase2_gf, i);
                                 auto [tid, il] = h2_hijack.match_name(t->name);
                                 if (tid >= 0) ggml_backend_sched_set_tensor_backend(sched.get(), t, gpu);
-                                if (t->op == GGML_OP_GLU) {
-                                    for (int s = 0; s < 2 && t->src[s]; s++)
-                                        ggml_backend_sched_set_tensor_backend(sched.get(), t->src[s], gpu);
-                                }
-                                if (t->op == GGML_OP_MUL_MAT_ID) {
-                                    for (int s = 0; s < GGML_MAX_SRC && t->src[s]; s++) {
-                                        if (s == 0) continue; // src[0]=weight: MoE cache GPU buffer, don't force
-                                        ggml_backend_sched_set_tensor_backend(sched.get(), t->src[s], gpu);
-                                    }
-                                }
-                            }
-                            for (int i = 0; i < ggml_graph_n_leafs(phase2_gf); i++) {
-                                ggml_tensor * t = ggml_graph_leaf(phase2_gf, i);
-                                auto [tid, il] = h2_hijack.match_name(t->name);
-                                if (tid >= 0) ggml_backend_sched_set_tensor_backend(sched.get(), t, gpu);
                             }
                             force_idxs_to_cpu();
                             if (!ggml_backend_sched_alloc_graph(sched.get(), phase2_gf)) {
@@ -2121,11 +2107,15 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll
                                 res->set_inputs(&ubatch);
                                 void* capture_st = ggml_backend_cuda_get_stream_ptr(gpu, 0);
                                 h2_hijack.copy_data_to_static(capture_st);
+                                fprintf(stderr, "phase2: D2D copy done, starting capture\n");
                                 ggml_backend_sched_synchronize(sched.get());
 
                                 // 5. CUDA Graph Capture
+                                fprintf(stderr, "phase2: beginCapture...\n");
                                 if (cudaStreamBeginCapture(capture_st, cudaStreamCaptureModeRelaxed) == cudaSuccess) {
+                                    fprintf(stderr, "phase2: capture started, computing...\n");
                                     auto s = graph_compute(phase2_gf, false);
+                                    fprintf(stderr, "phase2: graph_compute done status=%d\n", (int)s);
                                     if (s == GGML_STATUS_SUCCESS) {
                                         cudaGraph_t g;
                                         if (cudaStreamEndCapture(capture_st, &g) == cudaSuccess) {