Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/mmf.cuh (CUDA)

Code before:
#include "mma.cuh"
#include "common.cuh"
#include "convert.cuh"

using namespace ggml_cuda_mma;

#define MMF_ROWS_PER_BLOCK 32
#define MMF_ROWS_PER_BLOCK_CDNA 64

static __forceinline__ int64_t get_max_block_size(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 512;
    } else {
        return 256;
    }
}

static __forceinline__ int get_padded_size(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 2;
    } else {
        return 4;
    }
}

static constexpr __device__ int get_padded_size() {
#if defined(AMD_MFMA_AVAILABLE)
    return 2;
#else
    return 4;
#endif // defined(AMD_MFMA_AVAILABLE)
}

struct mmf_ids_data {
    const int32_t * ids_src_compact = nullptr;
    const int32_t * ids_dst_compact = nullptr;
    const int32_t * expert_bounds_dev = nullptr;
    int n_experts = 0;
    int sis1 = 0;
};

void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst);

bool ggml_cuda_should_use_mmf(enum ggml_type type, int cc, int warp_size, const int64_t * scr0_ne, const size_t * src0_nb, const int src1_ncols, bool mul_mat_id);

template <typename T, int rows_per_block, int cols_per_block, int nwarps, bool has_ids>
__launch_bounds__(ggml_cuda_get_physical_warp_size()*nwarps, 1)
static __global__ void mul_mat_f(
        const T * __restrict__ x, const float * __restrict__ y, const int32_t * __restrict__ ids, float * __restrict__ dst,
        const int ncols, const int ncols_dst_total, const int nchannels_dst, const int stride_row, const int stride_col_y, const int stride_col_dst,
        const int stride_col_id, const int stride_row_id,

Review comment (suggestion):
```suggestion
static constexpr __device__ int mmf_get_padding() {
```

Code after:
#include "mma.cuh"
#include "common.cuh"
#include "convert.cuh"

using namespace ggml_cuda_mma;

#define MMF_ROWS_PER_BLOCK 32
#define MMF_ROWS_PER_BLOCK_CDNA 64

static __forceinline__ int64_t mmf_get_max_block_size(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 512;
    } else {
        return 256;
    }
}

static __forceinline__ int mmf_get_padding(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 2;
    } else {
        return 4;
    }
}

static constexpr __device__ int mmf_get_padding() {
#if defined(AMD_MFMA_AVAILABLE)
    return 2;
#else
    return 4;
#endif // defined(AMD_MFMA_AVAILABLE)
}

struct mmf_ids_data {
    const int32_t * ids_src_compact = nullptr;
    const int32_t * ids_dst_compact = nullptr;
    const int32_t * expert_bounds_dev = nullptr;
    int n_experts = 0;
    int sis1 = 0;
};

void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst);

bool ggml_cuda_should_use_mmf(enum ggml_type type, int cc, int warp_size, const int64_t * scr0_ne, const size_t * src0_nb, const int src1_ncols, bool mul_mat_id);

template <typename T, int rows_per_block, int cols_per_block, int nwarps, bool has_ids>
__launch_bounds__(ggml_cuda_get_physical_warp_size()*nwarps, 1)
static __global__ void mul_mat_f(
        const T * __restrict__ x, const float * __restrict__ y, const int32_t * __restrict__ ids, float * __restrict__ dst,
        const int ncols, const int ncols_dst_total, const int nchannels_dst, const int stride_row, const int stride_col_y, const int stride_col_dst,
        const int stride_col_id, const int stride_row_id,