fix USE_CUDA_GRAPH=OFF

ggml-ci

Author: CISC

ggml/src/ggml-cuda/mean.cu

@@ -28,6 +28,7 @@ void ggml_cuda_op_mean(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
     cudaStreamCaptureStatus iscapturing;
     CUDA_CHECK(cudaStreamIsCapturing(stream, &iscapturing));
     if ((nrows == 1) &&
+#ifdef USE_CUDA_GRAPH
             // CUDA_GRAPHS_DISABLED
             ((ncols > 65536) &&
              ((ctx.cuda_graph->instance == nullptr) && (iscapturing == cudaStreamCaptureStatusNone) ||
@@ -38,6 +39,9 @@ void ggml_cuda_op_mean(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
          !((ctx.cuda_graph->instance == nullptr) && (iscapturing == cudaStreamCaptureStatusNone) ||
            ctx.cuda_graph->disable_due_to_gpu_arch || ctx.cuda_graph->disable_due_to_too_many_updates ||
            ctx.cuda_graph->disable_due_to_failed_graph_capture))) {
+#else
+        (ncols > 65536)) {
+#endif // USE_CUDA_GRAPH
         // Single row - use device-wide reduction
         size_t           tmp_size = 0;
         ggml_cuda_pool & pool     = ctx.pool();
@@ -51,7 +55,7 @@ void ggml_cuda_op_mean(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
         divide_by_count<float><<<1, 1, 0, stream>>>(dst_d, ncols);
         return;
     }
-#endif
+#endif // GGML_CUDA_USE_CUB
 
     const dim3 block_nums(nrows, 1, 1);