CUDA: always create events for split buffers

[JohannesGaessler]

Fixes #10176 .

I think the correct way to fix it is to just create the events unconditionally. Regardless of how the data is split you always need the events on the currently active device for the other devices to wait on. You could maybe reduce the number of events by only initializing those that are actually needed but I don't think that would be worthwhile since for the vast majority of use cases all events are already being created and used anyways.

[slaren]

Qwen2.5-0.5B does not work with this change alone, it still crashes in the memcpy later:

CUDA error: invalid argument
  current device: 1, in function ggml_cuda_op_mul_mat at ggml/src/ggml-cuda.cu:1583
  cudaMemcpyPeerAsync( src1_ddq_i, id, src1_ddq_i_source, ctx.device, src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs, stream)

[slaren]

It would also be possible to prevent using a split buffer entirely if the matrix is too small by returning false in the supports_op check.

[slaren]

This seems too expensive to do in this function, since this is called many times during inference by ggml_backend_sched. I think it should be possible to compute the minimum tensor size in ggml_backend_cuda_split_buffer_type, and store it in ggml_backend_cuda_split_buffer_type_context, then this function would only need to compare the tensor size to this value.