Divide mmf by ncols_dst

Author: am17an

ggml/src/ggml-cuda/mma.cuh

@@ -1,3 +1,4 @@
+#pragma once
 // This file contains primitives that expose the tensor core PTX instructions for CUDA code.
 // The primitives can be used in a similar way as the nvcuda::wmma interface but with a well-defined memory layout.
 // The documentation for the PTX instructions can be found under:

ggml/src/ggml-cuda/mmf.cu

@@ -163,7 +163,6 @@ static void launch_mul_mat_vec_f_cuda(
     const int nbytes_shared = warp_size*sizeof(float);
     const dim3 block_nums(nrows, nchannels_dst, nsamples_dst);
     const dim3 block_dims(block_size_best, 1, 1);
-
     switch (block_size_best) {
         case   32: {
             mul_mat_vec_f<T, type_acc, ncols_dst,  32><<<block_nums, block_dims, nbytes_shared, stream>>>

ggml/src/ggml-cuda/mmf.cuh

@@ -34,10 +34,6 @@
 DECL_MMQ_CASE({type});
 """
 
-TYPES_MMF = [
-    "GGML_TYPE_F32", "GGML_TYPE_F16", "GGML_TYPE_BF16"
-]
-
 SOURCE_MMF = """// This file has been autogenerated by generate_cu_files.py, do not edit manually.
 
 #include "../mmf.cuh"
@@ -88,6 +84,6 @@ def get_head_sizes(type_k, type_v):
     with open(f"mmq-instance-{get_short_name(type)}.cu", "w") as f:
         f.write(SOURCE_MMQ.format(type=type))
 
-for type in TYPES_MMF:
-    with open(f"mmf-instance-{get_short_name(type)}.cu", "w") as f:
+for type in range(1, 17):
+    with open(f"mmf-instance-ncols_{type}.cu", "w") as f:
         f.write(SOURCE_MMF.format(type=type))