diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu
index e9ffd274b..78ad607e3 100644
--- a/ggml/src/ggml-cuda/concat.cu
+++ b/ggml/src/ggml-cuda/concat.cu
@@ -123,33 +123,34 @@ static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE)
           uint64_t   nb0,
           uint64_t   nb1,
           uint64_t   nb2,
-          uint64_t   nb3){
+          uint64_t   nb3,
+               int   es){
     static_assert(dim >= 0 && dim <= 3, "dim must be in [0, 3]");
 
     const int64_t i3 = blockIdx.z;
     const int64_t i2 = blockIdx.y;
     const int64_t i1 = blockIdx.x;
 
-    const float * x;
+    const char * x;
 
     for (int64_t i0 = threadIdx.x; i0 < ne0; i0 += blockDim.x) {
         if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
-            x = (const float *)(src0 + (i3       )*nb03 + (i2       )*nb02 + (i1       )*nb01 + (i0       )*nb00);
+            x = (const char *)(src0 + (i3       )*nb03 + (i2       )*nb02 + (i1       )*nb01 + (i0       )*nb00);
         } else {
             if constexpr (dim == 0) {
-                x = (const float *) (src1 + i3 * nb13 + i2 * nb12 + i1 * nb11 + (i0 - ne00) * nb10);
+                x = (const char *) (src1 + i3 * nb13 + i2 * nb12 + i1 * nb11 + (i0 - ne00) * nb10);
             } else if constexpr (dim == 1) {
-                x = (const float *) (src1 + i3 * nb13 + i2 * nb12 + (i1 - ne01) * nb11 + i0 * nb10);
+                x = (const char *) (src1 + i3 * nb13 + i2 * nb12 + (i1 - ne01) * nb11 + i0 * nb10);
             } else if constexpr (dim == 2) {
-                x = (const float *) (src1 + i3 * nb13 + (i2 - ne02) * nb12 + i1 * nb11 + i0 * nb10);
+                x = (const char *) (src1 + i3 * nb13 + (i2 - ne02) * nb12 + i1 * nb11 + i0 * nb10);
             } else if constexpr (dim == 3) {
-                x = (const float *) (src1 + (i3 - ne03) * nb13 + i2 * nb12 + i1 * nb11 + i0 * nb10);
+                x = (const char *) (src1 + (i3 - ne03) * nb13 + i2 * nb12 + i1 * nb11 + i0 * nb10);
             }
         }
 
-        float * y = (float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
+        char * y = (char *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
 
-        *y = *x;
+        for (int b = 0; b < es; ++b) y[b] = x[b];   // element-size-generic copy (I32/F16/BF16/F32)
     }
 }
 
@@ -162,11 +163,13 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
 
     const int32_t dim = ((int32_t *) dst->op_params)[0];
 
-    GGML_ASSERT(src0->type == GGML_TYPE_F32);
-    GGML_ASSERT(src1->type == GGML_TYPE_F32);
-    GGML_ASSERT(dst->type  == GGML_TYPE_F32);
+    // DS4: the CSA path concats I32 (indices) and F16/BF16 (compressed KV), not just F32. Concat is
+    // pure placement (no arithmetic), so we copy raw element-size bytes. Require matching types;
+    // keep the F32/I32 (4-byte) contiguous fast path, route everything else through the generic kernel.
+    GGML_ASSERT(src1->type == src0->type);
+    GGML_ASSERT(dst->type  == src0->type);
 
-    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_type_size(dst->type) == 4) {
         const float * src0_d = (const float *)src0->data;
         const float * src1_d = (const float *)src1->data;
 
@@ -198,7 +201,8 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
                 src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
                 src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
                 dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
-                dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3]);
+                dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3],
+                (int) ggml_type_size(dst->type));
         };
         switch (dim) {
             case 0: