一、引言

1.1 Keye-VL-2.0-30B-A3B简介

快手Keye团队于2026年5月正式开源了Keye-VL-2.0-30B-A3B多模态大模型。该模型是Keye系列最新的30B级主力基座,首次将DSA(DeepSeek Sparse Attention)稀疏注意力机制引入多模态场景,支持256K超长上下文,实现了对小时级视频的毫秒级时序推理。模型采用MoE(混合专家)架构,总参数300亿,在多模态理解、长视频时序感知等任务上达到了同尺寸SOTA水平。快手官方同时提供了GGUF格式的量化模型,方便通过llama.cpp进行本地部署。

1.2 为什么在海光K100_AI上选择llama.cpp

在海光K100_AI(DCU)上部署大模型,目前面临严峻的生态困境:

(1)vLLM DCU适配滞后:国产芯片使用的vLLM版本通常较老,且对特定模型的支持需要逐一适配。海光提供的vLLM镜像往往只能运行特定版本的模型。
(2)SGLang支持有限:虽然SGLang正在向国产GPU扩展支持,但进展缓慢。
(3)Transformers兼容性存疑:PyTorch虽然通过CUDA API兼容层可以在海光DCU上运行,但多模态模型的复杂依赖往往存在各种兼容性问题。

在vLLM、SGLang、Transformers等主流推理框架都无法直接支持Keye-VL-2.0的情况下(需要大模型官方提供定制分支),llama.cpp凭借其轻量级、跨平台、对GGUF格式原生支持的特点,成为了在海光K100_AI上部署该模型的唯一可行选择。

1.3 环境说明

本文的部署环境如下:

项目规格
硬件海光K100_AI(8卡,gfx928架构)
操作系统Ubuntu 22.04
DTK版本DTK 26.04
目标模型Keye-VL-2.0-30B-A3B-BF16.gguf(约57GB)

二、部署测试详细步骤

2.1 模型下载

使用ModelScope下载Keye-VL-2.0-30B-A3B的GGUF模型文件:

# 下载主模型(BF16格式,约57GB)
modelscope download --model Kwai-Keye/Keye-VL-2.0-30B-A3B-GGUF \
    Keye-VL-2.0-30B-A3B-BF16.gguf \
    --local_dir /opt/models/Keye-VL-2.0-30B-A3B-GGUF

# 下载多模态投影文件(mmproj)
modelscope download --model Kwai-Keye/Keye-VL-2.0-30B-A3B-GGUF \
    mmproj-Keye-VL-2.0-30B-A3B-BF16.gguf \
    --local_dir /opt/models/Keye-VL-2.0-30B-A3B-GGUF

2.2 克隆快手定制版llama.cpp

快手官方提供了专门适配Keye-VL的llama.cpp定制分支:

git clone -b keye-vl-v2-30b-release-video https://github.com/Kwai-Keye/llama.cpp.git
cd llama.cpp

2.3 环境配置与代码修改

2.3.1 修改CMakeLists.txt适配DTK 26.04

由于DTK 26.04的某些版本检测逻辑与llama.cpp不兼容,需要修改CMake配置:

vim ggml/src/ggml-hip/CMakeLists.txt

找到VERSION_LESS 6.1,修改为:

VERSION_LESS 6.1 AND FALSE
2.3.2 修改DTK系统头文件(关键步骤)

海光DCU的gfx928架构硬限制线程块最大为256,但llama.cpp中大量GPU内核硬编码为1024线程。同时,DTK 26.04的cooperative_groups头文件与gfx928存在兼容性问题,需要手动修改系统头文件。

(1)修改_CG_ASM_PTR_CONSTRAINT

vim /opt/dtk-26.04/include/hip/amd_detail/cooperative_groups/details/info.h

cat /opt/dtk-26.04/include/hip/amd_detail/cooperative_groups/details/info.h | grep -B 5 -A 5 CG_ASM_PTR_CONSTRAINT
# define _CG_STATIC_CONST_DECL static const
# define _CG_CONST_DECL const
#endif

#if (defined(_MSC_VER) && !defined(_WIN64)) || defined(__arm__)
# define _CG_ASM_PTR_CONSTRAINT "r"
#else
#  define _CG_ASM_PTR_CONSTRAINT "r"
#endif

//__device_builtin__ void                   __trap(void);

# define _CG_ASSERT(x)


#将_CG_ASM_PTR_CONSTRAINT从"l"改为"r"(所有分支)。

(2)注释掉__any_sync函数

在同一文件中,找到__any_sync函数并将其注释掉,避免__any重复定义错误。

cat /opt/dtk-26.04/include/hip/amd_detail/cooperative_groups/details/info.h | grep -B 5 -A 5 __any_sync
}
#endif

//__device__
//inline
//int __any_sync(unsigned mask, int predicate){
//    return __any(predicate);
//}
#endif //__cplusplus
#endif // _CG_INFO_H_
2.3.3 修改llama.cpp源码适配gfx928

(1)强制使用MMQ内核

编辑ggml/src/ggml-cuda/mmf.cu,找到ggml_cuda_should_use_mmf函数,修改为:

bool ggml_cuda_should_use_mmf(enum ggml_type type, int cc, int warp_size, 
                              const int64_t * src0_ne, const size_t * src0_nb, 
                              const int src1_ncols, bool mul_mat_id) {
    return false;  // 强制禁用MFMA路径
}

(2)修改norm.cu中的所有1024为256

# 备份
cp ggml/src/ggml-cuda/norm.cu ggml/src/ggml-cuda/norm.cu.bak

# 将所有 if (ncols < 1024) 改为 if (ncols < 256)
sed -i 's/ncols < 1024/ncols < 256/g' ggml/src/ggml-cuda/norm.cu

# 将所有 block_dims(1024, 1, 1) 改为 block_dims(256, 1, 1)
sed -i 's/block_dims(1024, 1, 1)/block_dims(256, 1, 1)/g' ggml/src/ggml-cuda/norm.cu

# 将所有 norm_f32<1024> 改为 norm_f32<256>
sed -i 's/norm_f32<1024>/norm_f32<256>/g' ggml/src/ggml-cuda/norm.cu

# 将所有 group_norm_f32<1024> 改为 group_norm_f32<256>
sed -i 's/group_norm_f32<1024>/group_norm_f32<256>/g' ggml/src/ggml-cuda/norm.cu

# 将所有 rms_norm_back_f32<1024> 改为 rms_norm_back_f32<256>
sed -i 's/rms_norm_back_f32<1024>/rms_norm_back_f32<256>/g' ggml/src/ggml-cuda/norm.cu

# 将所有 l2_norm_f32<1024> 改为 l2_norm_f32<256>
sed -i 's/l2_norm_f32<1024>/l2_norm_f32<256>/g' ggml/src/ggml-cuda/norm.cu

# 修改 group_size 判断
sed -i 's/group_size < 1024/group_size < 256/g' ggml/src/ggml-cuda/norm.cu

(3)修改static_assert

编辑ggml/src/ggml-cuda/norm.cu,将第181行的静态断言修改为:

static_assert(block_size == 256, "unexpected block_size");

在文件开头添加宏定义:

#ifndef GGML_CUDA_MAX_THREADS_PER_BLOCK
#define GGML_CUDA_MAX_THREADS_PER_BLOCK 256
#endif

(4)修改softmax相关文件

# 检查并修改 softmax.cuh 中的宏
vim ggml/src/ggml-cuda/softmax.cuh
# 将 CUDA_SOFT_MAX_BLOCK_SIZE 从 1024 改为 256

# 修改 softmax.cu 中的静态断言
vim ggml/src/ggml-cuda/softmax.cu
# 第334行修改为:
# static_assert(CUDA_SOFT_MAX_BLOCK_SIZE == 256, "These values need to be adjusted.");

(5)修改common.cuh

vim ggml/src/ggml-cuda/common.cuh

#ifndef GGML_CUDA_MAX_THREADS_PER_BLOCK
#define GGML_CUDA_MAX_THREADS_PER_BLOCK 256
#endif

(6)修改其他文件中的block_dims

sed -i 's/block_dims(512,/block_dims(256,/g' ggml/src/ggml-cuda/mean.cu
sed -i 's/block_dims(512,/block_dims(256,/g' ggml/src/ggml-cuda/sumrows.cu

(7)修改./ggml/src/ggml-cuda/mmq.cuh的3951行代码为:

    const dim3 block_dims(warp_size, nwarps > 2 ? 2 : nwarps, 1);

(8)检查ggml/src/ggml-cuda/norm.cu代码修改情况,完整代码如下:

cat -n ggml/src/ggml-cuda/norm.cu
     1	#ifndef GGML_CUDA_MAX_THREADS_PER_BLOCK
     2	#define GGML_CUDA_MAX_THREADS_PER_BLOCK 256
     3	#endif
     4	#include "norm.cuh"
     5	#include <cstdint>
     6	
     7	template <int block_size>
     8	static __global__ void norm_f32(
     9	        const float * x, float * dst, const int ncols, const int64_t stride_row, const int64_t stride_channel,
    10	        const int64_t stride_sample, const float eps) {
    11	    const int nrows     = gridDim.x;
    12	    const int nchannels = gridDim.y;
    13	
    14	    const int row       = blockIdx.x;
    15	    const int channel   = blockIdx.y;
    16	    const int sample    = blockIdx.z;
    17	    const int tid       = threadIdx.x;
    18	
    19	    x   += sample*stride_sample + channel*stride_channel + row*stride_row;
    20	    dst += ((sample*nchannels + channel)*nrows + row)*ncols;
    21	
    22	    float2 mean_var = make_float2(0.0f, 0.0f);
    23	
    24	    ggml_cuda_pdl_sync();
    25	    for (int col = tid; col < ncols; col += block_size) {
    26	        const float xi = x[col];
    27	        mean_var.x += xi;
    28	        mean_var.y += xi * xi;
    29	    }
    30	
    31	    // sum up partial sums
    32	    extern __shared__ float2 s_sum2[];
    33	    mean_var = block_reduce<block_reduce_method::SUM, block_size>(mean_var, s_sum2);
    34	
    35	    const float mean = mean_var.x / ncols;
    36	    const float var = mean_var.y / ncols - mean * mean;
    37	    const float inv_std = rsqrtf(var + eps);
    38	
    39	    for (int col = tid; col < ncols; col += block_size) {
    40	        dst[col] = (x[col] - mean) * inv_std;
    41	    }
    42	}
    43	
    44	template <int block_size>
    45	static __global__ void group_norm_f32(const float * x, float * dst, const int group_size, const int ne_elements, const float eps) {
    46	    // blockIdx.x: num_groups idx
    47	    // threadIdx.x: block_size idx
    48	    const int start =     blockIdx.x*group_size + threadIdx.x;
    49	    const int end   = min(blockIdx.x*group_size + group_size,  ne_elements);
    50	
    51	    float tmp = 0.0f; // partial sum for thread in warp
    52	
    53	    ggml_cuda_pdl_sync();
    54	    for (int j = start; j < end; j += block_size) {
    55	        tmp += x[j];
    56	    }
    57	
    58	    extern __shared__ float s_sum[];
    59	    tmp = block_reduce<block_reduce_method::SUM, block_size>(tmp, s_sum);
    60	
    61	    const float mean = tmp / group_size;
    62	    tmp = 0.0f;
    63	
    64	    for (int j = start; j < end; j += block_size) {
    65	        const float xi = x[j] - mean;
    66	        dst[j] = xi;
    67	        tmp += xi * xi;
    68	    }
    69	
    70	    tmp = block_reduce<block_reduce_method::SUM, block_size>(tmp, s_sum);
    71	
    72	    const float variance = tmp / group_size;
    73	    const float scale = rsqrtf(variance + eps);
    74	    for (int j = start; j < end; j += block_size) {
    75	        dst[j] *= scale;
    76	    }
    77	}
    78	
    79	template <int block_size, bool do_multiply = false, bool do_add = false>
    80	__launch_bounds__(256, 1)
    81	static __global__ void rms_norm_f32(const float * x,
    82	                                    float *       dst,
    83	                                    const int     ncols,
    84	                                    const int64_t stride_row,
    85	                                    const int64_t stride_channel,
    86	                                    const int64_t stride_sample,
    87	                                    const float   eps,
    88	                                    const float * mul                  = nullptr,
    89	                                    const int64_t mul_stride_row       = 0,
    90	                                    const int64_t mul_stride_channel   = 0,
    91	                                    const int64_t mul_stride_sample    = 0,
    92	                                    const uint3   mul_ncols_packed     = make_uint3(0, 0, 0),
    93	                                    const uint3   mul_nrows_packed     = make_uint3(0, 0, 0),
    94	                                    const uint3   mul_nchannels_packed = make_uint3(0, 0, 0),
    95	                                    const uint3   mul_nsamples_packed  = make_uint3(0, 0, 0),
    96	                                    const float * add                  = nullptr,
    97	                                    const int64_t add_stride_row       = 0,
    98	                                    const int64_t add_stride_channel   = 0,
    99	                                    const int64_t add_stride_sample    = 0,
   100	                                    const uint3   add_ncols_packed     = make_uint3(0, 0, 0),
   101	                                    const uint3   add_nrows_packed     = make_uint3(0, 0, 0),
   102	                                    const uint3   add_nchannels_packed = make_uint3(0, 0, 0),
   103	                                    const uint3   add_nsamples_packed  = make_uint3(0, 0, 0)) {
   104	    ggml_cuda_pdl_lc();
   105	    const int nrows     = gridDim.x;
   106	    const int nchannels = gridDim.y;
   107	
   108	    const int row       = blockIdx.x;
   109	    const int channel   = blockIdx.y;
   110	    const int sample    = blockIdx.z;
   111	    const int tid       = threadIdx.x;
   112	
   113	    static_assert(!do_add || do_multiply, "fusing add is not supported without multiplying");
   114	
   115	    x   += sample*stride_sample + channel*stride_channel + row*stride_row;
   116	    dst += ((sample*nchannels + channel)*nrows + row)*ncols;
   117	
   118	    if constexpr (do_multiply) {
   119	        const uint32_t mul_row     = fastmodulo(row, mul_nrows_packed);
   120	        const uint32_t mul_channel = fastmodulo(channel, mul_nchannels_packed);
   121	        const uint32_t mul_sample  = fastmodulo(sample, mul_nsamples_packed);
   122	        mul += mul_sample * mul_stride_sample + mul_channel * mul_stride_channel + mul_row * mul_stride_row;
   123	    }
   124	
   125	    if constexpr (do_add) {
   126	        const int add_row     = fastmodulo(row, add_nrows_packed);
   127	        const int add_channel = fastmodulo(channel, add_nchannels_packed);
   128	        const int add_sample  = fastmodulo(sample, add_nsamples_packed);
   129	        add += add_sample * add_stride_sample + add_channel * add_stride_channel + add_row * add_stride_row;
   130	    }
   131	
   132	    float tmp = 0.0f; // partial sum for thread in warp
   133	
   134	    ggml_cuda_pdl_sync();
   135	    for (int col = tid; col < ncols; col += block_size) {
   136	        const float xi = x[col];
   137	        tmp += xi * xi;
   138	    }
   139	
   140	    // sum up partial sums
   141	    extern __shared__ float s_sum[];
   142	    tmp = block_reduce<block_reduce_method::SUM, block_size>(tmp, s_sum);
   143	
   144	    const float mean = tmp / ncols;
   145	    const float scale = rsqrtf(mean + eps);
   146	
   147	    for (int col = tid; col < ncols; col += block_size) {
   148	        if constexpr (do_multiply && do_add) {
   149	            const int mul_col = fastmodulo(col, mul_ncols_packed);
   150	            const int add_col = fastmodulo(col, add_ncols_packed);
   151	            dst[col]          = scale * x[col] * mul[mul_col] + add[add_col];
   152	        } else if constexpr (do_multiply) {
   153	            const int mul_col = fastmodulo(col, mul_ncols_packed);
   154	            dst[col]          = scale * x[col] * mul[mul_col];
   155	        } else {
   156	            dst[col] = scale * x[col];
   157	        }
   158	    }
   159	}
   160	
   161	template <int block_size>
   162	static __global__ void rms_norm_back_f32(
   163	        const float * grad, const float * xf, float * dst, const int ncols, const float eps) {
   164	    const int row = blockIdx.x*blockDim.y + threadIdx.y;
   165	    const int tid = threadIdx.x;
   166	
   167	    grad += int64_t(row)*ncols;
   168	    xf   += int64_t(row)*ncols;
   169	    dst  += int64_t(row)*ncols;
   170	
   171	    float sum_xx = 0.0f; // sum for squares of x, equivalent to forward pass
   172	    float sum_xg = 0.0f; // sum for x * gradient, needed because RMS norm mixes inputs
   173	
   174	    ggml_cuda_pdl_sync();
   175	    for (int col = tid; col < ncols; col += block_size) {
   176	        const float xfi = xf[col];
   177	        sum_xx += xfi * xfi;
   178	        sum_xg += xfi * grad[col];
   179	    }
   180	
   181	    // sum up partial sums
   182	    sum_xx = warp_reduce_sum(sum_xx);
   183	    sum_xg = warp_reduce_sum(sum_xg);
   184	    if constexpr (block_size > WARP_SIZE) {
   185	        static_assert(block_size == 256, "unexpected block_size");
   186	        __shared__ float s_sum_xx[32];
   187	        __shared__ float s_sum_xg[32];
   188	        const int warp_id = threadIdx.x / WARP_SIZE;
   189	        const int lane_id = threadIdx.x % WARP_SIZE;
   190	        if (lane_id == 0) {
   191	            s_sum_xx[warp_id] = sum_xx;
   192	            s_sum_xg[warp_id] = sum_xg;
   193	        }
   194	        __syncthreads();
   195	
   196	        sum_xx = s_sum_xx[lane_id];
   197	        sum_xx = warp_reduce_sum(sum_xx);
   198	
   199	        sum_xg = s_sum_xg[lane_id];
   200	        sum_xg = warp_reduce_sum(sum_xg);
   201	    }
   202	
   203	    const float mean_eps = sum_xx / ncols + eps;
   204	    const float sum_eps  = sum_xx + ncols*eps;
   205	
   206	    const float scale_grad = rsqrtf(mean_eps);
   207	    const float scale_x    = -scale_grad * sum_xg/sum_eps;
   208	
   209	    for (int col = tid; col < ncols; col += block_size) {
   210	        dst[col] = scale_grad*grad[col] + scale_x*xf[col];
   211	    }
   212	}
   213	
   214	// template <int block_size>
   215	// static __global__ void l2_norm_f32(const float * x, float * dst, const int ncols, const float eps) {
   216	//     const int row = blockIdx.x*blockDim.y + threadIdx.y;
   217	//     const int tid = threadIdx.x;
   218	
   219	//     float tmp = 0.0f; // partial sum for thread in warp
   220	
   221	//     for (int col = tid; col < ncols; col += block_size) {
   222	//         const float xi = x[row*ncols + col];
   223	//         tmp += xi * xi;
   224	//     }
   225	
   226	//     // sum up partial sums
   227	//     tmp = warp_reduce_sum(tmp);
   228	//     if (block_size > WARP_SIZE) {
   229	//         __shared__ float s_sum[32];
   230	//         int warp_id = threadIdx.x / WARP_SIZE;
   231	//         int lane_id = threadIdx.x % WARP_SIZE;
   232	//         if (lane_id == 0) {
   233	//             s_sum[warp_id] = tmp;
   234	//         }
   235	//         __syncthreads();
   236	//         tmp = s_sum[lane_id];
   237	//         tmp = warp_reduce_sum(tmp);
   238	//     }
   239	
   240	//     // from https://pytorch.org/docs/stable/generated/torch.nn.functional.normalize.html
   241	//     const float scale = rsqrtf(fmaxf(tmp, eps * eps));
   242	
   243	//     for (int col = tid; col < ncols; col += block_size) {
   244	//         dst[row*ncols + col] = scale * x[row*ncols + col];
   245	//     }
   246	// }
   247	
   248	template <int block_size>
   249	static __global__ void l2_norm_f32(
   250	        const float * x, float * dst, const int ncols, const int64_t stride_row, const int64_t stride_channel,
   251	        const int64_t stride_sample, const float eps) {
   252	    const int nrows     = gridDim.x;
   253	    const int nchannels = gridDim.y;
   254	
   255	    const int row       = blockIdx.x;
   256	    const int channel   = blockIdx.y;
   257	    const int sample    = blockIdx.z;
   258	    const int tid       = threadIdx.x;
   259	
   260	    x   += sample*stride_sample + channel*stride_channel + row*stride_row;
   261	    dst += ((sample*nchannels + channel)*nrows + row)*ncols;
   262	
   263	    float tmp = 0.0f; // partial sum for thread in warp
   264	
   265	    ggml_cuda_pdl_sync();
   266	    for (int col = tid; col < ncols; col += block_size) {
   267	        const float xi = x[col];
   268	        tmp += xi * xi;
   269	    }
   270	
   271	    // sum up partial sums
   272	    extern __shared__ float s_sum[];
   273	    tmp = block_reduce<block_reduce_method::SUM, block_size>(tmp, s_sum);
   274	    ggml_cuda_pdl_lc();
   275	
   276	    // from https://pytorch.org/docs/stable/generated/torch.nn.functional.normalize.html
   277	    const float scale = rsqrtf(fmaxf(tmp, eps * eps));
   278	
   279	    for (int col = tid; col < ncols; col += block_size) {
   280	        dst[col] = scale * x[col];
   281	    }
   282	}
   283	
   284	static void norm_f32_cuda(
   285	        const float * x, float * dst, const int ncols, const int nrows, const int nchannels, const int nsamples,
   286	        const int64_t stride_row, const int64_t stride_channel, const int64_t stride_sample, const float eps, cudaStream_t stream) {
   287	    const dim3 blocks_num(nrows, nchannels, nsamples);
   288	    if (ncols < 256) {
   289	        const dim3 block_dims(WARP_SIZE, 1, 1);
   290	        norm_f32<WARP_SIZE><<<blocks_num, block_dims, 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps);
   291	    } else {
   292	        const dim3 block_dims(256, 1, 1);
   293	        norm_f32<256><<<blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float2): 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps);
   294	    }
   295	}
   296	
   297	static void group_norm_f32_cuda(
   298	        const float * x, float * dst, const int num_groups, const float eps, const int group_size, const int ne_elements, cudaStream_t stream) {
   299	    if (group_size < 256) {
   300	        const dim3 block_dims(WARP_SIZE, 1, 1);
   301	        group_norm_f32<WARP_SIZE><<<num_groups, block_dims, 0, stream>>>(x, dst, group_size, ne_elements, eps);
   302	    } else {
   303	        const dim3 block_dims(256, 1, 1);
   304	        group_norm_f32<256><<<num_groups, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream>>>(x, dst, group_size, ne_elements, eps);
   305	    }
   306	}
   307	
   308	static void rms_norm_f32_cuda(
   309	        const float * x, float * dst, const int ncols, const int nrows, const int nchannels, const int nsamples,
   310	        const int64_t stride_row, const int64_t stride_channel, const int64_t stride_sample, const float eps, cudaStream_t stream) {
   311	    const dim3 blocks_num(nrows, nchannels, nsamples);
   312	    if (ncols < 256) {
   313	        const dim3 block_dims(256, 1, 1);
   314	        const ggml_cuda_kernel_launch_params launch_params = {blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   315	        ggml_cuda_kernel_launch(rms_norm_f32<256, false>, launch_params,
   316	            x, dst, ncols, stride_row, stride_channel, stride_sample, eps,
   317	        // underlying cudaLaunchKernelEx does not support default params
   318	        nullptr, 0, 0, 0, make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0),
   319	        nullptr, 0, 0, 0, make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0));
   320	    } else {
   321	        const dim3 block_dims(256, 1, 1);
   322	        const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   323	        ggml_cuda_kernel_launch(rms_norm_f32<256, false>, launch_params, x, dst, ncols, stride_row, stride_channel, stride_sample, eps,
   324	        // underlying cudaLaunchKernelEx does not support default params
   325	        nullptr, 0, 0, 0, make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0),
   326	        nullptr, 0, 0, 0, make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0));
   327	    }
   328	}
   329	
   330	static void rms_norm_mul_f32_cuda(const float *  x,
   331	                                  const float *  mul,
   332	                                  const float *  add,
   333	                                  float *        dst,
   334	                                  const int      ncols,
   335	                                  const int      nrows,
   336	                                  const int      nchannels,
   337	                                  const int      nsamples,
   338	                                  const int64_t  stride_row,
   339	                                  const int64_t  stride_channel,
   340	                                  const int64_t  stride_sample,
   341	                                  const int64_t  mul_stride_row,
   342	                                  const int64_t  mul_stride_channel,
   343	                                  const int64_t  mul_stride_sample,
   344	                                  const uint32_t mul_ncols,
   345	                                  const uint32_t mul_nrows,
   346	                                  const uint32_t mul_nchannels,
   347	                                  const uint32_t mul_nsamples,
   348	                                  const int64_t  add_stride_row,
   349	                                  const int64_t  add_stride_channel,
   350	                                  const int64_t  add_stride_sample,
   351	                                  const uint32_t add_ncols,
   352	                                  const uint32_t add_nrows,
   353	                                  const uint32_t add_nchannels,
   354	                                  const uint32_t add_nsamples,
   355	                                  const float    eps,
   356	                                  cudaStream_t   stream) {
   357	    const dim3 blocks_num(nrows, nchannels, nsamples);
   358	    if (mul == nullptr) {
   359	        rms_norm_f32_cuda(x, dst, ncols, nrows, nchannels, nsamples, stride_row, stride_channel, stride_sample, eps, stream);
   360	        return;
   361	    }
   362	    if (add == nullptr) {
   363	        const uint3 mul_ncols_packed     = init_fastdiv_values(mul_ncols);
   364	        const uint3 mul_nrows_packed     = init_fastdiv_values(mul_nrows);
   365	        const uint3 mul_nchannels_packed = init_fastdiv_values(mul_nchannels);
   366	        const uint3 mul_nsamples_packed  = init_fastdiv_values(mul_nsamples);
   367	        if (ncols < 256) {
   368	            const dim3 block_dims(256, 1, 1);
   369	            const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   370	            ggml_cuda_kernel_launch(rms_norm_f32<256, true>, launch_params,
   371	                x, dst, ncols, stride_row, stride_channel, stride_sample, eps, mul, mul_stride_row, mul_stride_channel,
   372	                mul_stride_sample, mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
   373	                // underlying cudaLaunchKernelEx does not support default params
   374	            nullptr, 0, 0, 0, make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0));
   375	        } else {
   376	            const dim3 block_dims(256, 1, 1);
   377	            const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   378	            ggml_cuda_kernel_launch(rms_norm_f32<256, true>, launch_params,
   379	                x, dst, ncols, stride_row, stride_channel, stride_sample, eps, mul, mul_stride_row, mul_stride_channel,
   380	                mul_stride_sample, mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed,
   381	                // underlying cudaLaunchKernelEx does not support default params
   382	            nullptr, 0, 0, 0, make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0), make_uint3(0, 0, 0));
   383	        }
   384	    } else {
   385	        const uint3 mul_ncols_packed     = init_fastdiv_values(mul_ncols);
   386	        const uint3 mul_nrows_packed     = init_fastdiv_values(mul_nrows);
   387	        const uint3 mul_nchannels_packed = init_fastdiv_values(mul_nchannels);
   388	        const uint3 mul_nsamples_packed  = init_fastdiv_values(mul_nsamples);
   389	
   390	        const uint3 add_ncols_packed     = init_fastdiv_values(add_ncols);
   391	        const uint3 add_nrows_packed     = init_fastdiv_values(add_nrows);
   392	        const uint3 add_nchannels_packed = init_fastdiv_values(add_nchannels);
   393	        const uint3 add_nsamples_packed  = init_fastdiv_values(add_nsamples);
   394	        if (ncols < 256) {
   395	            const dim3 block_dims(256, 1, 1);
   396	            const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims,block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   397	            ggml_cuda_kernel_launch(rms_norm_f32<256, true, true>, launch_params,
   398	                x, dst, ncols, stride_row, stride_channel, stride_sample, eps, mul, mul_stride_row, mul_stride_channel,
   399	                mul_stride_sample, mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, add,
   400	                add_stride_row, add_stride_channel, add_stride_sample, add_ncols_packed, add_nrows_packed,
   401	                add_nchannels_packed, add_nsamples_packed);
   402	        } else {
   403	            const dim3 block_dims(256, 1, 1);
   404	            const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   405	            ggml_cuda_kernel_launch(rms_norm_f32<256, true, true>, launch_params,
   406	                x, dst, ncols, stride_row, stride_channel, stride_sample, eps, mul, mul_stride_row, mul_stride_channel,
   407	                mul_stride_sample, mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, add,
   408	                add_stride_row, add_stride_channel, add_stride_sample, add_ncols_packed, add_nrows_packed,
   409	                add_nchannels_packed, add_nsamples_packed);
   410	        }
   411	    }
   412	}
   413	
   414	static void rms_norm_back_f32_cuda(const float * grad, const float * xf, float * dst, const int ncols, const int nrows, const float eps, cudaStream_t stream) {
   415	    if (ncols < 256) {
   416	        const dim3 block_dims(WARP_SIZE, 1, 1);
   417	        rms_norm_back_f32<WARP_SIZE><<<nrows, block_dims, 0, stream>>>(grad, xf, dst, ncols, eps);
   418	    } else {
   419	        const dim3 block_dims(256, 1, 1);
   420	        rms_norm_back_f32<256><<<nrows, block_dims, 0, stream>>>(grad, xf, dst, ncols, eps);
   421	    }
   422	}
   423	
   424	static void l2_norm_f32_cuda(
   425	        const float * x, float * dst, const int ncols, const int nrows, const int nchannels, const int nsamples,
   426	        const int64_t stride_row, const int64_t stride_channel, const int64_t stride_sample, const float eps, cudaStream_t stream) {
   427	    const dim3 blocks_num(nrows, nchannels, nsamples);
   428	    if (ncols < 256) {
   429	        const dim3 block_dims(WARP_SIZE, 1, 1);
   430	        const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims, 0, stream};
   431	        ggml_cuda_kernel_launch(l2_norm_f32<WARP_SIZE>, launch_params, x, dst, ncols, stride_row, stride_channel, stride_sample, eps);
   432	    } else {
   433	        const dim3 block_dims(256, 1, 1);
   434	        const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params{blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream};
   435	        ggml_cuda_kernel_launch(l2_norm_f32<256>, launch_params, x, dst, ncols, stride_row, stride_channel, stride_sample, eps);
   436	    }
   437	}
   438	
   439	void ggml_cuda_op_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
   440	    const ggml_tensor * src0 = dst->src[0];
   441	    const float * src0_d = (const float *) src0->data;
   442	    float * dst_d = (float *) dst->data;
   443	    cudaStream_t stream = ctx.stream();
   444	
   445	    GGML_ASSERT(src0->type == GGML_TYPE_F32);
   446	    GGML_ASSERT( dst->type == GGML_TYPE_F32);
   447	
   448	    GGML_TENSOR_UNARY_OP_LOCALS;
   449	
   450	    float eps;
   451	    memcpy(&eps, dst->op_params, sizeof(float));
   452	    GGML_ASSERT(eps >= 0.0f);
   453	
   454	    const size_t ts0 = ggml_type_size(src0->type);
   455	    GGML_ASSERT(nb00 == ts0);
   456	    const int64_t s01 = nb01 / ts0;
   457	    const int64_t s02 = nb02 / ts0;
   458	    const int64_t s03 = nb03 / ts0;
   459	
   460	    norm_f32_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s01, s02, s03, eps, stream);
   461	}
   462	
   463	void ggml_cuda_op_group_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
   464	    const ggml_tensor * src0 = dst->src[0];
   465	    const float * src0_d = (const float *)src0->data;
   466	    float * dst_d = (float *)dst->data;
   467	    cudaStream_t stream = ctx.stream();
   468	
   469	    GGML_ASSERT(src0->type == GGML_TYPE_F32);
   470	    GGML_ASSERT( dst->type == GGML_TYPE_F32);
   471	
   472	    int num_groups = dst->op_params[0];
   473	
   474	    float eps;
   475	    memcpy(&eps, dst->op_params + 1, sizeof(float));
   476	    GGML_ASSERT(eps >= 0.0f);
   477	
   478	    int group_size = src0->ne[0] * src0->ne[1] * ((src0->ne[2] + num_groups - 1) / num_groups);
   479	    group_norm_f32_cuda(src0_d, dst_d, num_groups * src0->ne[3], eps, group_size, ggml_nelements(src0), stream);
   480	}
   481	
   482	void ggml_cuda_op_rms_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
   483	    const ggml_tensor * src0 = dst->src[0];
   484	    const float * src0_d = (const float *) src0->data;
   485	    float * dst_d = (float *) dst->data;
   486	    cudaStream_t stream = ctx.stream();
   487	
   488	    GGML_ASSERT(src0->type == GGML_TYPE_F32);
   489	    GGML_ASSERT( dst->type == GGML_TYPE_F32);
   490	
   491	    GGML_TENSOR_UNARY_OP_LOCALS;
   492	
   493	    float eps;
   494	    memcpy(&eps, dst->op_params, sizeof(float));
   495	    GGML_ASSERT(eps >= 0.0f);
   496	
   497	    const size_t ts0 = ggml_type_size(src0->type);
   498	    GGML_ASSERT(nb00 == ts0);
   499	    const int64_t s01 = nb01 / ts0;
   500	    const int64_t s02 = nb02 / ts0;
   501	    const int64_t s03 = nb03 / ts0;
   502	
   503	    rms_norm_f32_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s01, s02, s03, eps, stream);
   504	}
   505	
   506	void ggml_cuda_op_rms_norm_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * mul_tensor) {
   507	    const ggml_tensor * rms_norm_src = (ggml_tensor *) dst->src[0];
   508	    float eps = 0.0f;
   509	
   510	    memcpy(&eps, dst->op_params, sizeof(float));
   511	
   512	    const float * src0_d = (const float *) rms_norm_src->data;
   513	    const float * mul_d = nullptr;
   514	    const ggml_tensor * mul_src = nullptr;
   515	
   516	    if (mul_tensor->src[0] == dst) {
   517	        mul_d = (float *) mul_tensor->src[1]->data;
   518	        mul_src = mul_tensor->src[1];
   519	    } else if(mul_tensor->src[1] == dst) {
   520	        mul_d = (float *) mul_tensor->src[0]->data;
   521	        mul_src = mul_tensor->src[0];
   522	    } else {
   523	        GGML_ASSERT(false);
   524	    }
   525	
   526	    float * dst_d = (float *) mul_tensor->data;
   527	    cudaStream_t stream = ctx.stream();
   528	
   529	    GGML_ASSERT(rms_norm_src->type == GGML_TYPE_F32);
   530	    GGML_ASSERT(dst->type == GGML_TYPE_F32);
   531	    GGML_ASSERT(mul_tensor->type == GGML_TYPE_F32);
   532	    GGML_ASSERT(eps >= 0.0f);
   533	
   534	    const int64_t ne00 = rms_norm_src->ne[0];
   535	    const int64_t ne01 = rms_norm_src->ne[1];
   536	    const int64_t ne02 = rms_norm_src->ne[2];
   537	    const int64_t ne03 = rms_norm_src->ne[3];
   538	
   539	    const size_t ts0 = ggml_type_size(rms_norm_src->type);
   540	    GGML_ASSERT(rms_norm_src->nb[0] == ts0);
   541	    const int64_t s01 = rms_norm_src->nb[1] / ts0;
   542	    const int64_t s02 = rms_norm_src->nb[2] / ts0;
   543	    const int64_t s03 = rms_norm_src->nb[3] / ts0;
   544	
   545	    const size_t ts_mul = ggml_type_size(mul_src->type);
   546	    GGML_ASSERT(mul_src->nb[0] == ts_mul);
   547	    const int64_t mul_s01 = mul_src->nb[1] / ts_mul;
   548	    const int64_t mul_s02 = mul_src->nb[2] / ts_mul;
   549	    const int64_t mul_s03 = mul_src->nb[3] / ts_mul;
   550	
   551	    const int mul_ncols     = mul_src->ne[0];
   552	    const int mul_nrows     = mul_src->ne[1];
   553	    const int mul_nchannels = mul_src->ne[2];
   554	    const int mul_nsamples  = mul_src->ne[3];
   555	
   556	    rms_norm_mul_f32_cuda(src0_d, mul_d, nullptr, dst_d,
   557	                          ne00, ne01, ne02, ne03,
   558	                          /*s00*/ s01, s02, s03,
   559	                          /*mul_s00*/ mul_s01, mul_s02, mul_s03,
   560	                          mul_ncols, mul_nrows, mul_nchannels, mul_nsamples,
   561	                          /*add_s00*/ 0, 0, 0,
   562	                          0, 0, 0, 0,
   563	                          eps, stream);
   564	}
   565	
   566	void ggml_cuda_op_rms_norm_fused_add(ggml_backend_cuda_context & ctx,
   567	                                     ggml_tensor *               dst,
   568	                                     ggml_tensor *               mul_tensor,
   569	                                     ggml_tensor *               add_tensor) {
   570	    const ggml_tensor * rms_norm_src = (ggml_tensor *) dst->src[0];
   571	    float               eps          = 0.0f;
   572	
   573	    memcpy(&eps, dst->op_params, sizeof(float));
   574	
   575	    const float *       src0_d  = (const float *) rms_norm_src->data;
   576	    const float *       mul_d   = nullptr;
   577	    const ggml_tensor * mul_src = nullptr;
   578	
   579	    if (mul_tensor->src[0] == dst) {
   580	        mul_d   = (float *) mul_tensor->src[1]->data;
   581	        mul_src = mul_tensor->src[1];
   582	    } else if (mul_tensor->src[1] == dst) {
   583	        mul_d   = (float *) mul_tensor->src[0]->data;
   584	        mul_src = mul_tensor->src[0];
   585	    } else {
   586	        GGML_ASSERT(false);
   587	    }
   588	
   589	    const float *       add_d   = nullptr;
   590	    const ggml_tensor * add_src = nullptr;
   591	
   592	    if (add_tensor->src[0] == mul_tensor) {
   593	        add_d   = (float *) add_tensor->src[1]->data;
   594	        add_src = add_tensor->src[1];
   595	    } else if (add_tensor->src[1] == mul_tensor) {
   596	        add_d   = (float *) add_tensor->src[0]->data;
   597	        add_src = add_tensor->src[0];
   598	    } else {
   599	        GGML_ASSERT(false);
   600	    }
   601	
   602	    float *      dst_d  = (float *) add_tensor->data;
   603	    cudaStream_t stream = ctx.stream();
   604	
   605	    GGML_ASSERT(rms_norm_src->type == GGML_TYPE_F32);
   606	    GGML_ASSERT(dst->type == GGML_TYPE_F32);
   607	    GGML_ASSERT(mul_tensor->type == GGML_TYPE_F32);
   608	    GGML_ASSERT(add_tensor->type == GGML_TYPE_F32);
   609	    GGML_ASSERT(eps >= 0.0f);
   610	
   611	    const int64_t ne00 = rms_norm_src->ne[0];
   612	    const int64_t ne01 = rms_norm_src->ne[1];
   613	    const int64_t ne02 = rms_norm_src->ne[2];
   614	    const int64_t ne03 = rms_norm_src->ne[3];
   615	
   616	    const size_t ts0 = ggml_type_size(rms_norm_src->type);
   617	    GGML_ASSERT(rms_norm_src->nb[0] == ts0);
   618	    const int64_t s01 = rms_norm_src->nb[1] / ts0;
   619	    const int64_t s02 = rms_norm_src->nb[2] / ts0;
   620	    const int64_t s03 = rms_norm_src->nb[3] / ts0;
   621	
   622	    const size_t ts_mul = ggml_type_size(mul_src->type);
   623	    GGML_ASSERT(mul_src->nb[0] == ts_mul);
   624	    const int64_t mul_s01 = mul_src->nb[1] / ts_mul;
   625	    const int64_t mul_s02 = mul_src->nb[2] / ts_mul;
   626	    const int64_t mul_s03 = mul_src->nb[3] / ts_mul;
   627	
   628	    const int mul_ncols     = mul_src->ne[0];
   629	    const int mul_nrows     = mul_src->ne[1];
   630	    const int mul_nchannels = mul_src->ne[2];
   631	    const int mul_nsamples  = mul_src->ne[3];
   632	
   633	    const size_t ts_add = ggml_type_size(add_src->type);
   634	    GGML_ASSERT(add_src->nb[0] == ts_add);
   635	    const int64_t add_s01 = add_src->nb[1] / ts_add;
   636	    const int64_t add_s02 = add_src->nb[2] / ts_add;
   637	    const int64_t add_s03 = add_src->nb[3] / ts_add;
   638	
   639	    const int add_ncols     = add_src->ne[0];
   640	    const int add_nrows     = add_src->ne[1];
   641	    const int add_nchannels = add_src->ne[2];
   642	    const int add_nsamples  = add_src->ne[3];
   643	
   644	    rms_norm_mul_f32_cuda(src0_d, mul_d,add_d,dst_d,
   645	                          ne00,ne01, ne02, ne03,
   646	                          /*s00*/ s01, s02, s03,
   647	                          /*mul_s00*/ mul_s01, mul_s02, mul_s03,
   648	                          mul_ncols, mul_nrows, mul_nchannels, mul_nsamples,
   649	                          /*add_s00*/ add_s01, add_s02, add_s03,
   650	                          add_ncols, add_nrows, add_nchannels, add_nsamples,
   651	                          eps, stream);
   652	}
   653	
   654	void ggml_cuda_op_rms_norm_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
   655	    const ggml_tensor * grad  = dst->src[0]; // gradients
   656	    const ggml_tensor * src0f = dst->src[1]; // src0 from forward pass
   657	
   658	    const float * grad_d  = (const float *) grad->data;
   659	    const float * src0f_d = (const float *) src0f->data;
   660	    float       * dst_d   = (float       *) dst->data;
   661	
   662	    cudaStream_t stream = ctx.stream();
   663	
   664	    GGML_ASSERT(ggml_is_contiguous(grad));
   665	
   666	    GGML_ASSERT( grad->type == GGML_TYPE_F32);
   667	    GGML_ASSERT(src0f->type == GGML_TYPE_F32);
   668	    GGML_ASSERT(  dst->type == GGML_TYPE_F32);
   669	
   670	    const int64_t ne00 = src0f->ne[0];
   671	    const int64_t nrows = ggml_nrows(src0f);
   672	
   673	    float eps;
   674	    memcpy(&eps, dst->op_params, sizeof(float));
   675	    GGML_ASSERT(eps >= 0.0f);
   676	
   677	    rms_norm_back_f32_cuda(grad_d, src0f_d, dst_d, ne00, nrows, eps, stream);
   678	}
   679	
   680	void ggml_cuda_op_l2_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
   681	    const ggml_tensor * src0 = dst->src[0];
   682	    const float * src0_d = (const float *) src0->data;
   683	    float * dst_d = (float *) dst->data;
   684	    cudaStream_t stream = ctx.stream();
   685	
   686	    GGML_ASSERT(src0->type == GGML_TYPE_F32);
   687	    GGML_ASSERT( dst->type == GGML_TYPE_F32);
   688	
   689	    GGML_TENSOR_UNARY_OP_LOCALS;
   690	
   691	    float eps;
   692	    memcpy(&eps, dst->op_params, sizeof(float));
   693	    GGML_ASSERT(eps >= 0.0f);
   694	
   695	    const size_t ts0 = ggml_type_size(src0->type);
   696	    GGML_ASSERT(nb00 == ts0);
   697	    const int64_t s01 = nb01 / ts0;
   698	    const int64_t s02 = nb02 / ts0;
   699	    const int64_t s03 = nb03 / ts0;
   700	
   701	    l2_norm_f32_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s01, s02, s03, eps, stream);
   702	}
root@:/opt/models/Kwai-llama.cpp-new# grep -n "1024" ggml/src/ggml-cuda/norm.cu | grep -v "48 \* 1024"
root@:/opt/models/Kwai-llama.cpp-new# 
root@:/opt/models/Kwai-llama.cpp-new# grep -n "256" ggml/src/ggml-cuda/norm.cu 
2:#define GGML_CUDA_MAX_THREADS_PER_BLOCK 256
80:__launch_bounds__(256, 1)
185:        static_assert(block_size == 256, "unexpected block_size");
288:    if (ncols < 256) {
292:        const dim3 block_dims(256, 1, 1);
293:        norm_f32<256><<<blocks_num, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float2): 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps);
299:    if (group_size < 256) {
303:        const dim3 block_dims(256, 1, 1);
304:        group_norm_f32<256><<<num_groups, block_dims, block_dims.x > WARP_SIZE ? 32 * sizeof(float): 0, stream>>>(x, dst, group_size, ne_elements, eps);
312:    if (ncols < 256) {
313:        const dim3 block_dims(256, 1, 1);
315:        ggml_cuda_kernel_launch(rms_norm_f32<256, false>, launch_params,
321:        const dim3 block_dims(256, 1, 1);
323:        ggml_cuda_kernel_launch(rms_norm_f32<256, false>, launch_params, x, dst, ncols, stride_row, stride_channel, stride_sample, eps,
367:        if (ncols < 256) {
368:            const dim3 block_dims(256, 1, 1);
370:            ggml_cuda_kernel_launch(rms_norm_f32<256, true>, launch_params,
376:            const dim3 block_dims(256, 1, 1);
378:            ggml_cuda_kernel_launch(rms_norm_f32<256, true>, launch_params,
394:        if (ncols < 256) {
395:            const dim3 block_dims(256, 1, 1);
397:            ggml_cuda_kernel_launch(rms_norm_f32<256, true, true>, launch_params,
403:            const dim3 block_dims(256, 1, 1);
405:            ggml_cuda_kernel_launch(rms_norm_f32<256, true, true>, launch_params,
415:    if (ncols < 256) {
419:        const dim3 block_dims(256, 1, 1);
420:        rms_norm_back_f32<256><<<nrows, block_dims, 0, stream>>>(grad, xf, dst, ncols, eps);
428:    if (ncols < 256) {
433:        const dim3 block_dims(256, 1, 1);
435:        ggml_cuda_kernel_launch(l2_norm_f32<256>, launch_params, x, dst, ncols, stride_row, stride_channel, stride_sample, eps);

2.4 编译

HIPCXX="$(hipconfig -l)/clang" \
HIP_PATH="$(hipconfig -R)" \
cmake -S . -B build \
    -DGGML_HIP=ON \
    -DGPU_TARGETS=gfx928 \
    -DCMAKE_BUILD_TYPE=Release \
    -DGGML_CUDA_FORCE_MMQ=ON

cmake --build build --config Release -j $(nproc) 2>&1 | tee build.log

2.5 启动服务

./build/bin/llama-server \
    -m /opt/models/Keye-VL-2.0-30B-A3B-GGUF/Keye-VL-2.0-30B-A3B-BF16.gguf \
    --mmproj /opt/models/Keye-VL-2.0-30B-A3B-GGUF/mmproj-Keye-VL-2.0-30B-A3B-BF16.gguf \
    --host 0.0.0.0 \
    --port 8082 \
    --alias Keye-VL-2.0-30B-A3B \
    --flash-attn off

注意:由于gfx928上Flash Attention的MFMA内核缺失,必须添加--flash-attn off参数禁用Flash Attention,否则服务会崩溃。

2.6 测试程序代码

cat test-video.py 
import json
import requests
import base64
import cv2
import os
import sys

BASE_URL = "http://192.168.222.65:8082"
MODEL_NAME = "Keye-VL-2.0-30B-A3B"   # 与 --alias 一致

def extract_frames(video_path, fps=2.0, max_frames=10):
    """
    从视频中按指定 fps 采样,返回帧的 Base64 列表(JPEG)。
    """
    cap = cv2.VideoCapture(video_path)
    if not cap.isOpened():
        raise RuntimeError(f"无法打开视频: {video_path}")
    video_fps = cap.get(cv2.CAP_PROP_FPS)
    total_frames = int(cap.get(cv2.CAP_PROP_FRAME_COUNT))
    if total_frames == 0:
        cap.release()
        raise RuntimeError("视频无有效帧")

    # 计算采样间隔(帧数)
    interval = max(1, int(video_fps / fps))
    frame_indices = list(range(0, total_frames, interval))
    if max_frames and len(frame_indices) > max_frames:
        # 均匀采样至 max_frames 张
        step = len(frame_indices) // max_frames
        frame_indices = frame_indices[::step][:max_frames]

    frames_b64 = []
    for idx in frame_indices:
        cap.set(cv2.CAP_PROP_POS_FRAMES, idx)
        ret, frame = cap.read()
        if not ret:
            continue
        _, buffer = cv2.imencode('.jpg', frame)
        img_b64 = base64.b64encode(buffer).decode('utf-8')
        frames_b64.append(img_b64)
    cap.release()
    return frames_b64

def generate(messages):
    payload = {
        "model": MODEL_NAME,
        "messages": messages,
        "n": 1,
        "temperature": 0.1,
        "max_tokens": 32760,
        "top_p": 0.001,
        "ignore_eos": False,
        "skip_special_tokens": True,
    }
    resp = requests.post(
        f"{BASE_URL}/v1/chat/completions",
        headers={"Content-Type": "application/json"},
        data=json.dumps(payload),
        timeout=1800,
    )
    resp.raise_for_status()
    return resp.json()

if __name__ == "__main__":
    video_path = "/opt/models/llama.cpp/Demo-乌海吵架.mp4"
    if not os.path.exists(video_path):
        print(f"视频不存在: {video_path}")
        sys.exit(1)

    # 抽取帧(按 2 fps,最多 10 张,避免 payload 过大)
    try:
        frames = extract_frames(video_path, fps=2.0, max_frames=10)
    except Exception as e:
        print(f"抽帧失败: {e}")
        sys.exit(1)

    if not frames:
        print("未能提取任何帧")
        sys.exit(1)

    # 构建 content:每张帧作为 image_url,最后添加文本
    content = []
    for b64 in frames:
        content.append({
            "type": "image_url",
            "image_url": {"url": f"data:image/jpeg;base64,{b64}"}
        })
    content.append({"type": "text", "text": "请详细描述这个视频的内容,结合所有帧理解场景和动作。"})

    messages = [{"role": "user", "content": content}]

    # 调用模型
    try:
        result = generate(messages)
        print("模型回答:")
        print(result["choices"][0]["message"]["content"])
    except Exception as e:
        print(f"请求失败: {e}")

2.7 大模型运行日志

./build/bin/llama-server     -m /opt/models/Keye-VL-2.0-30B-A3B-GGUF/Keye-VL-2.0-30B-A3B-BF16.gguf     --mmproj /opt/models/Keye-VL-2.0-30B-A3B-GGUF/mmproj-Keye-VL-2.0-30B-A3B-BF16.gguf     --host 0.0.0.0 --port 8082     --alias Keye-VL-2.0-30B-A3B --flash-attn off
0.00.268.210 I log_info: verbosity = 3 (adjust with the `-lv N` CLI arg)
0.00.268.217 I device_info:
0.00.268.267 I   - ROCm0   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.278 I   - ROCm1   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.283 I   - ROCm2   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.288 I   - ROCm3   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.292 I   - ROCm4   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.297 I   - ROCm5   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.302 I   - ROCm6   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.307 I   - ROCm7   : K100_AI (65520 MiB, 64686 MiB free)
0.00.268.321 I   - CPU     : Hygon C86-4G (OPN:7470) (515686 MiB, 515686 MiB free)
0.00.268.448 I system_info: n_threads = 96 (n_threads_batch = 96) / 192 | ROCm : FORCE_MMQ = 1 | NO_VMM = 1 | PEER_MAX_BATCH_SIZE = 128 | CPU : SSE3 = 1 | SSSE3 = 1 | AVX = 1 | AVX2 = 1 | F16C = 1 | FMA = 1 | BMI2 = 1 | LLAMAFILE = 1 | OPENMP = 1 | REPACK = 1 | 
0.00.268.458 I srv  llama_server: n_parallel is set to auto, using n_parallel = 4 and kv_unified = true
0.00.268.540 I srv          init: running without SSL
0.00.268.641 I srv          init: using 191 threads for HTTP server
0.00.269.080 I srv         start: binding port with default address family
0.00.270.455 I srv  llama_server: loading model
0.00.270.543 I srv    load_model: loading model '/opt/models/Keye-VL-2.0-30B-A3B-GGUF/Keye-VL-2.0-30B-A3B-BF16.gguf'
0.00.392.070 I srv    load_model: [mtmd] estimated worst-case memory usage of mmproj is 1518.16 MiB (took 121.47 ms)
0.00.392.124 I common_init_result: fitting params to device memory ...
0.00.392.129 I common_init_result: (for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)
0.00.984.652 W load: control-looking token: 128247 '</s>' was not control-type; this is probably a bug in the model. its type will be overridden
1.53.597.137 I common_init_from_params: warming up the model with an empty run - please wait ... (--no-warmup to disable)
2.09.040.449 I srv    load_model: loaded multimodal model, '/opt/models/Keye-VL-2.0-30B-A3B-GGUF/mmproj-Keye-VL-2.0-30B-A3B-BF16.gguf'
2.09.040.486 I srv    load_model: initializing slots, n_slots = 4
2.09.282.742 W common_speculative_init: no implementations specified for speculative decoding
2.09.282.763 I slot   load_model: id  0 | task -1 | new slot, n_ctx = 262144
2.09.282.783 I slot   load_model: id  1 | task -1 | new slot, n_ctx = 262144
2.09.282.785 I slot   load_model: id  2 | task -1 | new slot, n_ctx = 262144
2.09.282.787 I slot   load_model: id  3 | task -1 | new slot, n_ctx = 262144
2.09.283.038 I srv    load_model: prompt cache is enabled, size limit: 8192 MiB
2.09.283.045 I srv    load_model: use `--cache-ram 0` to disable the prompt cache
2.09.283.047 I srv    load_model: for more info see https://github.com/ggml-org/llama.cpp/pull/16391
2.09.283.048 I srv    load_model: context checkpoints enabled, max = 32, min spacing = 256
2.09.283.079 I srv          init: idle slots will be saved to prompt cache and cleared upon starting a new task
2.09.296.940 I init: chat template, example_format: '<|im_start|>system
You are a helpful assistant<|im_end|>
<|im_start|>user
Hello<|im_end|>
<|im_start|>assistant
Hi there<|im_end|>
<|im_start|>user
How are you?<|im_end|>
<|im_start|>assistant
'
2.09.305.215 I srv          init: init: chat template, thinking = 0
2.09.305.309 I srv  llama_server: model loaded
2.09.305.324 I srv  llama_server: server is listening on http://0.0.0.0:8082
2.09.305.334 I srv  update_slots: all slots are idle
2.17.222.807 I srv    operator(): Chat format: peg-native
2.17.223.015 I slot get_availabl: id  3 | task -1 | selected slot by LRU, t_last = -1
2.17.223.021 I srv  get_availabl: updating prompt cache
2.17.223.037 I srv          load:  - looking for better prompt, base f_keep = -1.000, sim = 0.000
2.17.223.053 I srv        update:  - cache state: 0 prompts, 0.000 MiB (limits: 8192.000 MiB, 262144 tokens, 8589934592 est)
2.17.223.056 I srv  get_availabl: prompt cache update took 0.03 ms
2.17.223.186 I slot launch_slot_: id  3 | task 0 | processing task, is_child = 0
2.17.223.193 I slot process_sing: id  0 | task -1 | saving idle slot to prompt cache
2.17.223.196 I slot prompt_clear: id  0 | task -1 | clearing prompt with 0 tokens
2.17.224.047 I slot process_sing: id  1 | task -1 | saving idle slot to prompt cache
2.17.224.053 I slot prompt_clear: id  1 | task -1 | clearing prompt with 0 tokens
2.17.224.713 I slot process_sing: id  2 | task -1 | saving idle slot to prompt cache
2.17.224.719 I slot prompt_clear: id  2 | task -1 | clearing prompt with 0 tokens
2.25.109.083 I slot print_timing: id  3 | task 0 | n_decoded =    100, tg =  18.12 t/s, tg_3s =  18.12 t/s
2.27.731.900 I slot print_timing: id  3 | task 0 | prompt eval time =    2365.68 ms /    10 tokens (  236.57 ms per token,     4.23 tokens per second)
2.27.731.912 I slot print_timing: id  3 | task 0 |        eval time =    8140.72 ms /   147 tokens (   55.38 ms per token,    18.06 tokens per second)
2.27.731.917 I slot print_timing: id  3 | task 0 |       total time =   10506.40 ms /   157 tokens
2.27.731.927 I slot print_timing: id  3 | task 0 |    graphs reused =          0
2.27.731.993 I slot      release: id  3 | task 0 | stop processing: n_tokens = 156, truncated = 0
2.27.732.020 I srv  update_slots: all slots are idle
2.38.259.866 I srv    operator(): Chat format: peg-native
2.38.260.043 I slot get_availabl: id  3 | task -1 | selected slot by LCP similarity, sim_best = 1.000 (> 0.100 thold), f_keep = 0.064
2.38.260.050 I srv  get_availabl: updating prompt cache
2.38.261.122 W srv   prompt_save:  - saving prompt with length 156, total state size = 17.374 MiB (draft: 0.000 MiB)
2.38.629.156 I srv          load:  - looking for better prompt, base f_keep = 0.064, sim = 1.000
2.38.629.173 I srv        update:  - cache state: 1 prompts, 17.374 MiB (limits: 8192.000 MiB, 262144 tokens, 262144 est)
2.38.629.176 I srv        update:    - prompt 0x56564dc74e00:     156 tokens, checkpoints:  0,    17.374 MiB
2.38.629.179 I srv  get_availabl: prompt cache update took 369.13 ms
2.38.629.330 I slot launch_slot_: id  3 | task 148 | processing task, is_child = 0
2.38.629.335 I slot process_sing: id  0 | task -1 | saving idle slot to prompt cache
2.38.629.336 I slot prompt_clear: id  0 | task -1 | clearing prompt with 0 tokens
2.38.629.699 I slot process_sing: id  1 | task -1 | saving idle slot to prompt cache
2.38.629.703 I slot prompt_clear: id  1 | task -1 | clearing prompt with 0 tokens
2.38.630.057 I slot process_sing: id  2 | task -1 | saving idle slot to prompt cache
2.38.630.061 I slot prompt_clear: id  2 | task -1 | clearing prompt with 0 tokens
2.38.630.464 W slot   operator(): id  3 | task 148 | need to evaluate at least 1 token for each active slot (n_past = 10, task.n_tokens() = 10)
2.38.630.468 W slot   operator(): id  3 | task 148 | n_past was set to 9
2.44.171.472 I slot print_timing: id  3 | task 148 | n_decoded =    100, tg =  18.22 t/s, tg_3s =  18.22 t/s
2.46.571.096 I slot print_timing: id  3 | task 148 | prompt eval time =      53.37 ms /     1 tokens (   53.37 ms per token,    18.74 tokens per second)
2.46.571.109 I slot print_timing: id  3 | task 148 |        eval time =    7887.22 ms /   143 tokens (   55.16 ms per token,    18.13 tokens per second)
2.46.571.112 I slot print_timing: id  3 | task 148 |       total time =    7940.58 ms /   144 tokens
2.46.571.115 I slot print_timing: id  3 | task 148 |    graphs reused =          0
2.46.571.151 I slot      release: id  3 | task 148 | stop processing: n_tokens = 152, truncated = 0
2.46.571.171 I srv  update_slots: all slots are idle
3.24.338.342 I srv    operator(): Chat format: peg-native
3.24.338.578 I slot get_availabl: id  2 | task -1 | selected slot by LRU, t_last = -1
3.24.338.586 I srv  get_availabl: updating prompt cache
3.24.338.600 I srv          load:  - looking for better prompt, base f_keep = -1.000, sim = 0.000
3.24.338.611 I srv        update:  - cache state: 1 prompts, 17.374 MiB (limits: 8192.000 MiB, 262144 tokens, 262144 est)
3.24.338.613 I srv        update:    - prompt 0x56564dc74e00:     156 tokens, checkpoints:  0,    17.374 MiB
3.24.338.615 I srv  get_availabl: prompt cache update took 0.03 ms
3.24.338.761 I slot launch_slot_: id  2 | task 292 | processing task, is_child = 0
3.24.338.766 I slot process_sing: id  0 | task -1 | saving idle slot to prompt cache
3.24.338.767 I slot prompt_clear: id  0 | task -1 | clearing prompt with 0 tokens
3.24.339.236 I slot process_sing: id  1 | task -1 | saving idle slot to prompt cache
3.24.339.241 I slot prompt_clear: id  1 | task -1 | clearing prompt with 0 tokens
3.24.339.595 I slot process_sing: id  3 | task -1 | saving idle slot to prompt cache
3.24.340.130 W srv   prompt_save:  - saving prompt with length 152, total state size = 16.929 MiB (draft: 0.000 MiB)
3.24.686.253 I srv        update:  - cache state: 2 prompts, 34.303 MiB (limits: 8192.000 MiB, 262144 tokens, 262144 est)
3.24.686.266 I srv        update:    - prompt 0x56564dc74e00:     156 tokens, checkpoints:  0,    17.374 MiB
3.24.686.267 I srv        update:    - prompt 0x56564dc9f6c0:     152 tokens, checkpoints:  0,    16.929 MiB
3.24.686.269 I slot prompt_clear: id  3 | task -1 | clearing prompt with 152 tokens
3.24.777.659 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 5, n_chunks = 1
3.29.552.528 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   1202, progress = 0.10, t =   4.87 s / 247.03 tokens per second
3.29.553.929 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 1202, n_chunks = 1
3.35.933.353 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   2399, progress = 0.20, t =  11.25 s / 213.31 tokens per second
3.35.934.685 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 2399, n_chunks = 1
3.46.409.143 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   3596, progress = 0.30, t =  21.72 s / 165.54 tokens per second
3.46.410.478 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 3596, n_chunks = 1
3.55.255.645 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   4793, progress = 0.40, t =  30.57 s / 156.79 tokens per second
3.55.256.991 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 4793, n_chunks = 1
4.06.581.606 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   5990, progress = 0.50, t =  41.89 s / 142.98 tokens per second
4.06.582.941 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 5990, n_chunks = 1
4.24.143.973 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   7187, progress = 0.60, t =  59.46 s / 120.88 tokens per second
4.24.145.305 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 7187, n_chunks = 1
4.36.096.484 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   8384, progress = 0.70, t =  71.41 s / 117.41 tokens per second
4.36.097.353 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 8384, n_chunks = 1
4.49.310.707 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =   9581, progress = 0.80, t =  84.62 s / 113.22 tokens per second
4.49.311.537 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 9581, n_chunks = 1
5.29.609.012 I slot print_timing: id  2 | task 292 | prompt processing, n_tokens =  10778, progress = 0.90, t = 124.92 s / 86.28 tokens per second
5.29.609.748 I slot process_mtmd: id  2 | task 292 | encoding mtmd batch from idx = 10778, n_chunks = 1
5.55.543.599 I srv    operator(): Chat format: peg-native
5.55.730.587 I slot get_availabl: id  1 | task -1 | selected slot by LRU, t_last = -1
5.55.730.597 I srv  get_availabl: updating prompt cache
5.55.730.611 I srv          load:  - looking for better prompt, base f_keep = -1.000, sim = 0.000
5.55.730.627 I srv        update:  - cache state: 2 prompts, 34.303 MiB (limits: 8192.000 MiB, 262144 tokens, 262144 est)
5.55.730.631 I srv        update:    - prompt 0x56564dc74e00:     156 tokens, checkpoints:  0,    17.374 MiB
5.55.730.638 I srv        update:    - prompt 0x56564dc9f6c0:     152 tokens, checkpoints:  0,    16.929 MiB
5.55.730.640 I srv  get_availabl: prompt cache update took 0.04 ms
5.55.730.813 I slot launch_slot_: id  1 | task 343 | processing task, is_child = 0
5.55.730.820 I slot process_sing: id  0 | task -1 | saving idle slot to prompt cache
5.55.730.822 I slot prompt_clear: id  0 | task -1 | clearing prompt with 0 tokens
5.55.731.860 I slot process_sing: id  3 | task -1 | saving idle slot to prompt cache
5.55.731.867 I slot prompt_clear: id  3 | task -1 | clearing prompt with 0 tokens
6.15.791.082 I slot print_timing: id  2 | task 292 | n_decoded =    100, tg =   3.27 t/s, tg_3s =   3.27 t/s
6.18.970.176 I slot print_timing: id  2 | task 292 | n_decoded =    111, tg =   3.29 t/s, tg_3s =   3.46 t/s
6.22.220.326 W srv          stop: cancel task, id_task = 343
6.22.220.872 I slot print_timing: id  2 | task 292 | n_decoded =    122, tg =   3.30 t/s, tg_3s =   3.38 t/s
6.22.220.901 I slot      release: id  1 | task 343 | stop processing: n_tokens = 199, truncated = 0
6.25.309.108 I slot print_timing: id  2 | task 292 | n_decoded =    133, tg =   3.32 t/s, tg_3s =   3.56 t/s
6.28.380.477 I slot print_timing: id  2 | task 292 | n_decoded =    144, tg =   3.34 t/s, tg_3s =   3.58 t/s
6.31.489.056 I slot print_timing: id  2 | task 292 | n_decoded =    155, tg =   3.35 t/s, tg_3s =   3.54 t/s
6.34.561.220 I slot print_timing: id  2 | task 292 | n_decoded =    166, tg =   3.37 t/s, tg_3s =   3.58 t/s
6.37.602.923 I slot print_timing: id  2 | task 292 | n_decoded =    177, tg =   3.38 t/s, tg_3s =   3.62 t/s
6.40.652.696 I slot print_timing: id  2 | task 292 | n_decoded =    188, tg =   3.39 t/s, tg_3s =   3.61 t/s
6.43.728.431 I slot print_timing: id  2 | task 292 | n_decoded =    199, tg =   3.40 t/s, tg_3s =   3.58 t/s
6.46.804.177 I slot print_timing: id  2 | task 292 | n_decoded =    210, tg =   3.41 t/s, tg_3s =   3.58 t/s
6.49.870.539 I slot print_timing: id  2 | task 292 | n_decoded =    221, tg =   3.42 t/s, tg_3s =   3.59 t/s
6.52.951.230 I slot print_timing: id  2 | task 292 | n_decoded =    232, tg =   3.43 t/s, tg_3s =   3.57 t/s
6.56.035.512 I slot print_timing: id  2 | task 292 | n_decoded =    243, tg =   3.43 t/s, tg_3s =   3.57 t/s
6.59.138.559 I slot print_timing: id  2 | task 292 | n_decoded =    254, tg =   3.44 t/s, tg_3s =   3.54 t/s
7.02.225.594 I slot print_timing: id  2 | task 292 | n_decoded =    265, tg =   3.44 t/s, tg_3s =   3.56 t/s
7.05.331.091 I slot print_timing: id  2 | task 292 | n_decoded =    276, tg =   3.45 t/s, tg_3s =   3.54 t/s
7.08.437.668 I slot print_timing: id  2 | task 292 | n_decoded =    287, tg =   3.45 t/s, tg_3s =   3.54 t/s
7.11.559.274 I slot print_timing: id  2 | task 292 | n_decoded =    298, tg =   3.45 t/s, tg_3s =   3.52 t/s
7.14.684.591 I slot print_timing: id  2 | task 292 | n_decoded =    309, tg =   3.46 t/s, tg_3s =   3.52 t/s
7.17.844.130 I slot print_timing: id  2 | task 292 | n_decoded =    320, tg =   3.46 t/s, tg_3s =   3.48 t/s
7.20.988.930 I slot print_timing: id  2 | task 292 | n_decoded =    331, tg =   3.46 t/s, tg_3s =   3.50 t/s
7.24.049.457 I slot print_timing: id  2 | task 292 | n_decoded =    342, tg =   3.46 t/s, tg_3s =   3.59 t/s
7.27.141.945 I slot print_timing: id  2 | task 292 | n_decoded =    353, tg =   3.46 t/s, tg_3s =   3.56 t/s
7.30.244.295 I slot print_timing: id  2 | task 292 | n_decoded =    364, tg =   3.47 t/s, tg_3s =   3.55 t/s
7.33.339.926 I slot print_timing: id  2 | task 292 | n_decoded =    375, tg =   3.47 t/s, tg_3s =   3.55 t/s
7.36.461.116 I slot print_timing: id  2 | task 292 | n_decoded =    386, tg =   3.47 t/s, tg_3s =   3.52 t/s
7.39.581.003 I slot print_timing: id  2 | task 292 | n_decoded =    397, tg =   3.47 t/s, tg_3s =   3.53 t/s
7.42.684.348 I slot print_timing: id  2 | task 292 | n_decoded =    408, tg =   3.47 t/s, tg_3s =   3.54 t/s
7.45.785.152 I slot print_timing: id  2 | task 292 | n_decoded =    419, tg =   3.48 t/s, tg_3s =   3.55 t/s
7.48.882.844 I slot print_timing: id  2 | task 292 | n_decoded =    430, tg =   3.48 t/s, tg_3s =   3.55 t/s
7.51.997.768 I slot print_timing: id  2 | task 292 | n_decoded =    441, tg =   3.48 t/s, tg_3s =   3.53 t/s
7.55.123.414 I slot print_timing: id  2 | task 292 | n_decoded =    452, tg =   3.48 t/s, tg_3s =   3.52 t/s
7.58.226.086 I slot print_timing: id  2 | task 292 | n_decoded =    463, tg =   3.48 t/s, tg_3s =   3.55 t/s
8.01.367.108 I slot print_timing: id  2 | task 292 | n_decoded =    474, tg =   3.48 t/s, tg_3s =   3.50 t/s
8.04.456.070 I slot print_timing: id  2 | task 292 | n_decoded =    485, tg =   3.48 t/s, tg_3s =   3.56 t/s
8.07.541.386 I slot print_timing: id  2 | task 292 | n_decoded =    496, tg =   3.49 t/s, tg_3s =   3.57 t/s
8.10.608.195 I slot print_timing: id  2 | task 292 | n_decoded =    507, tg =   3.49 t/s, tg_3s =   3.59 t/s
8.13.686.503 I slot print_timing: id  2 | task 292 | n_decoded =    518, tg =   3.49 t/s, tg_3s =   3.57 t/s
8.16.761.426 I slot print_timing: id  2 | task 292 | n_decoded =    529, tg =   3.49 t/s, tg_3s =   3.58 t/s
8.19.839.138 I slot print_timing: id  2 | task 292 | n_decoded =    540, tg =   3.49 t/s, tg_3s =   3.57 t/s
8.22.915.473 I slot print_timing: id  2 | task 292 | n_decoded =    551, tg =   3.49 t/s, tg_3s =   3.58 t/s
8.25.996.889 I slot print_timing: id  2 | task 292 | n_decoded =    562, tg =   3.50 t/s, tg_3s =   3.57 t/s
8.29.150.986 I slot print_timing: id  2 | task 292 | n_decoded =    573, tg =   3.50 t/s, tg_3s =   3.49 t/s
8.32.282.655 I slot print_timing: id  2 | task 292 | n_decoded =    584, tg =   3.50 t/s, tg_3s =   3.51 t/s
8.35.434.844 I slot print_timing: id  2 | task 292 | n_decoded =    595, tg =   3.50 t/s, tg_3s =   3.49 t/s
8.38.620.898 I slot print_timing: id  2 | task 292 | n_decoded =    606, tg =   3.50 t/s, tg_3s =   3.45 t/s
8.41.813.834 I slot print_timing: id  2 | task 292 | n_decoded =    617, tg =   3.49 t/s, tg_3s =   3.45 t/s
8.45.040.669 I slot print_timing: id  2 | task 292 | n_decoded =    628, tg =   3.49 t/s, tg_3s =   3.41 t/s
8.48.176.104 I slot print_timing: id  2 | task 292 | n_decoded =    639, tg =   3.49 t/s, tg_3s =   3.51 t/s
8.51.282.549 I slot print_timing: id  2 | task 292 | n_decoded =    650, tg =   3.49 t/s, tg_3s =   3.54 t/s
8.54.416.792 I slot print_timing: id  2 | task 292 | n_decoded =    661, tg =   3.49 t/s, tg_3s =   3.51 t/s
8.57.560.081 I slot print_timing: id  2 | task 292 | n_decoded =    672, tg =   3.49 t/s, tg_3s =   3.50 t/s
9.00.687.084 I slot print_timing: id  2 | task 292 | n_decoded =    683, tg =   3.49 t/s, tg_3s =   3.52 t/s
9.03.795.136 I slot print_timing: id  2 | task 292 | n_decoded =    694, tg =   3.50 t/s, tg_3s =   3.54 t/s
9.06.919.265 I slot print_timing: id  2 | task 292 | n_decoded =    705, tg =   3.50 t/s, tg_3s =   3.52 t/s
9.10.057.182 I slot print_timing: id  2 | task 292 | n_decoded =    716, tg =   3.50 t/s, tg_3s =   3.51 t/s
9.13.171.094 I slot print_timing: id  2 | task 292 | n_decoded =    727, tg =   3.50 t/s, tg_3s =   3.53 t/s
9.16.288.462 I slot print_timing: id  2 | task 292 | n_decoded =    738, tg =   3.50 t/s, tg_3s =   3.53 t/s
9.19.413.066 I slot print_timing: id  2 | task 292 | n_decoded =    749, tg =   3.50 t/s, tg_3s =   3.52 t/s
9.22.502.167 I slot print_timing: id  2 | task 292 | n_decoded =    760, tg =   3.50 t/s, tg_3s =   3.56 t/s
9.25.581.441 I slot print_timing: id  2 | task 292 | n_decoded =    771, tg =   3.50 t/s, tg_3s =   3.57 t/s
9.28.669.079 I slot print_timing: id  2 | task 292 | n_decoded =    782, tg =   3.50 t/s, tg_3s =   3.56 t/s
9.31.762.030 I slot print_timing: id  2 | task 292 | n_decoded =    793, tg =   3.50 t/s, tg_3s =   3.56 t/s
9.34.857.483 I slot print_timing: id  2 | task 292 | n_decoded =    804, tg =   3.50 t/s, tg_3s =   3.55 t/s
9.37.977.695 I slot print_timing: id  2 | task 292 | n_decoded =    815, tg =   3.50 t/s, tg_3s =   3.53 t/s
9.41.067.093 I slot print_timing: id  2 | task 292 | n_decoded =    826, tg =   3.50 t/s, tg_3s =   3.56 t/s
9.44.185.321 I slot print_timing: id  2 | task 292 | n_decoded =    837, tg =   3.50 t/s, tg_3s =   3.53 t/s
9.47.311.641 I slot print_timing: id  2 | task 292 | n_decoded =    848, tg =   3.50 t/s, tg_3s =   3.52 t/s
9.50.430.014 I slot print_timing: id  2 | task 292 | n_decoded =    859, tg =   3.50 t/s, tg_3s =   3.53 t/s
9.53.572.784 I slot print_timing: id  2 | task 292 | n_decoded =    870, tg =   3.50 t/s, tg_3s =   3.50 t/s
9.56.733.178 I slot print_timing: id  2 | task 292 | n_decoded =    881, tg =   3.50 t/s, tg_3s =   3.48 t/s
9.59.884.941 I slot print_timing: id  2 | task 292 | n_decoded =    892, tg =   3.50 t/s, tg_3s =   3.49 t/s
10.03.096.274 I slot print_timing: id  2 | task 292 | n_decoded =    903, tg =   3.50 t/s, tg_3s =   3.43 t/s
10.06.332.485 I slot print_timing: id  2 | task 292 | n_decoded =    914, tg =   3.50 t/s, tg_3s =   3.40 t/s
10.09.561.224 I slot print_timing: id  2 | task 292 | n_decoded =    925, tg =   3.50 t/s, tg_3s =   3.41 t/s
10.12.730.407 I slot print_timing: id  2 | task 292 | n_decoded =    936, tg =   3.50 t/s, tg_3s =   3.47 t/s
10.15.870.252 I slot print_timing: id  2 | task 292 | n_decoded =    947, tg =   3.50 t/s, tg_3s =   3.50 t/s
10.19.026.813 I slot print_timing: id  2 | task 292 | n_decoded =    958, tg =   3.50 t/s, tg_3s =   3.48 t/s
10.22.163.140 I slot print_timing: id  2 | task 292 | n_decoded =    969, tg =   3.50 t/s, tg_3s =   3.51 t/s
10.25.293.008 I slot print_timing: id  2 | task 292 | n_decoded =    980, tg =   3.50 t/s, tg_3s =   3.51 t/s
10.28.440.258 I slot print_timing: id  2 | task 292 | n_decoded =    991, tg =   3.50 t/s, tg_3s =   3.50 t/s
10.31.588.729 I slot print_timing: id  2 | task 292 | n_decoded =   1002, tg =   3.50 t/s, tg_3s =   3.49 t/s
10.34.731.825 I slot print_timing: id  2 | task 292 | n_decoded =   1013, tg =   3.50 t/s, tg_3s =   3.50 t/s
10.37.879.502 I slot print_timing: id  2 | task 292 | n_decoded =   1024, tg =   3.50 t/s, tg_3s =   3.49 t/s
10.41.029.036 I slot print_timing: id  2 | task 292 | n_decoded =   1035, tg =   3.50 t/s, tg_3s =   3.49 t/s
10.44.042.233 I slot print_timing: id  2 | task 292 | n_decoded =   1046, tg =   3.50 t/s, tg_3s =   3.65 t/s
10.47.228.108 I slot print_timing: id  2 | task 292 | n_decoded =   1057, tg =   3.50 t/s, tg_3s =   3.45 t/s
10.50.402.142 I slot print_timing: id  2 | task 292 | n_decoded =   1068, tg =   3.50 t/s, tg_3s =   3.47 t/s
10.53.565.582 I slot print_timing: id  2 | task 292 | n_decoded =   1079, tg =   3.50 t/s, tg_3s =   3.48 t/s
10.56.743.552 I slot print_timing: id  2 | task 292 | n_decoded =   1090, tg =   3.50 t/s, tg_3s =   3.46 t/s
10.59.927.867 I slot print_timing: id  2 | task 292 | n_decoded =   1101, tg =   3.50 t/s, tg_3s =   3.45 t/s
11.03.140.914 I slot print_timing: id  2 | task 292 | n_decoded =   1112, tg =   3.50 t/s, tg_3s =   3.42 t/s
11.06.343.984 I slot print_timing: id  2 | task 292 | n_decoded =   1123, tg =   3.50 t/s, tg_3s =   3.43 t/s
11.09.561.444 I slot print_timing: id  2 | task 292 | n_decoded =   1134, tg =   3.50 t/s, tg_3s =   3.42 t/s
11.12.816.549 I slot print_timing: id  2 | task 292 | n_decoded =   1145, tg =   3.50 t/s, tg_3s =   3.38 t/s
11.16.095.096 I slot print_timing: id  2 | task 292 | n_decoded =   1156, tg =   3.49 t/s, tg_3s =   3.36 t/s
11.19.390.999 I slot print_timing: id  2 | task 292 | n_decoded =   1167, tg =   3.49 t/s, tg_3s =   3.34 t/s
11.22.669.575 I slot print_timing: id  2 | task 292 | n_decoded =   1178, tg =   3.49 t/s, tg_3s =   3.36 t/s
11.25.860.787 I slot print_timing: id  2 | task 292 | n_decoded =   1189, tg =   3.49 t/s, tg_3s =   3.45 t/s
11.29.146.274 I slot print_timing: id  2 | task 292 | n_decoded =   1200, tg =   3.49 t/s, tg_3s =   3.35 t/s
11.32.351.665 I slot print_timing: id  2 | task 292 | n_decoded =   1211, tg =   3.49 t/s, tg_3s =   3.43 t/s
11.35.547.278 I slot print_timing: id  2 | task 292 | n_decoded =   1222, tg =   3.49 t/s, tg_3s =   3.44 t/s
11.38.765.224 I slot print_timing: id  2 | task 292 | n_decoded =   1233, tg =   3.49 t/s, tg_3s =   3.42 t/s
11.41.977.728 I slot print_timing: id  2 | task 292 | n_decoded =   1244, tg =   3.49 t/s, tg_3s =   3.42 t/s
11.45.174.870 I slot print_timing: id  2 | task 292 | n_decoded =   1255, tg =   3.49 t/s, tg_3s =   3.44 t/s
11.48.380.764 I slot print_timing: id  2 | task 292 | n_decoded =   1266, tg =   3.49 t/s, tg_3s =   3.43 t/s
11.51.612.185 I slot print_timing: id  2 | task 292 | n_decoded =   1277, tg =   3.49 t/s, tg_3s =   3.40 t/s
11.54.845.680 I slot print_timing: id  2 | task 292 | n_decoded =   1288, tg =   3.48 t/s, tg_3s =   3.40 t/s
11.58.066.148 I slot print_timing: id  2 | task 292 | n_decoded =   1299, tg =   3.48 t/s, tg_3s =   3.42 t/s
12.01.298.788 I slot print_timing: id  2 | task 292 | n_decoded =   1310, tg =   3.48 t/s, tg_3s =   3.40 t/s
12.04.529.448 I slot print_timing: id  2 | task 292 | n_decoded =   1321, tg =   3.48 t/s, tg_3s =   3.40 t/s
12.07.795.319 I slot print_timing: id  2 | task 292 | n_decoded =   1332, tg =   3.48 t/s, tg_3s =   3.37 t/s
12.10.851.131 I slot print_timing: id  2 | task 292 | n_decoded =   1342, tg =   3.48 t/s, tg_3s =   3.27 t/s
12.14.121.999 I slot print_timing: id  2 | task 292 | n_decoded =   1353, tg =   3.48 t/s, tg_3s =   3.36 t/s
12.17.370.012 I slot print_timing: id  2 | task 292 | n_decoded =   1364, tg =   3.48 t/s, tg_3s =   3.39 t/s
12.20.634.197 I slot print_timing: id  2 | task 292 | n_decoded =   1375, tg =   3.48 t/s, tg_3s =   3.37 t/s
12.23.661.224 I slot print_timing: id  2 | task 292 | n_decoded =   1385, tg =   3.48 t/s, tg_3s =   3.30 t/s
12.26.675.588 I slot print_timing: id  2 | task 292 | n_decoded =   1395, tg =   3.48 t/s, tg_3s =   3.32 t/s
12.29.968.772 I slot print_timing: id  2 | task 292 | n_decoded =   1406, tg =   3.47 t/s, tg_3s =   3.34 t/s
12.33.258.493 I slot print_timing: id  2 | task 292 | n_decoded =   1417, tg =   3.47 t/s, tg_3s =   3.34 t/s
12.36.301.918 I slot print_timing: id  2 | task 292 | n_decoded =   1427, tg =   3.47 t/s, tg_3s =   3.29 t/s
12.39.314.341 I slot print_timing: id  2 | task 292 | n_decoded =   1437, tg =   3.47 t/s, tg_3s =   3.32 t/s
12.42.354.874 I slot print_timing: id  2 | task 292 | n_decoded =   1447, tg =   3.47 t/s, tg_3s =   3.29 t/s
12.45.362.139 I slot print_timing: id  2 | task 292 | n_decoded =   1457, tg =   3.47 t/s, tg_3s =   3.33 t/s
12.48.375.484 I slot print_timing: id  2 | task 292 | n_decoded =   1467, tg =   3.47 t/s, tg_3s =   3.32 t/s
12.51.451.246 I slot print_timing: id  2 | task 292 | n_decoded =   1477, tg =   3.47 t/s, tg_3s =   3.25 t/s
12.54.523.943 I slot print_timing: id  2 | task 292 | n_decoded =   1487, tg =   3.46 t/s, tg_3s =   3.25 t/s
12.57.598.120 I slot print_timing: id  2 | task 292 | n_decoded =   1497, tg =   3.46 t/s, tg_3s =   3.25 t/s
13.00.699.093 I slot print_timing: id  2 | task 292 | n_decoded =   1507, tg =   3.46 t/s, tg_3s =   3.22 t/s
13.03.738.671 I slot print_timing: id  2 | task 292 | n_decoded =   1517, tg =   3.46 t/s, tg_3s =   3.29 t/s
13.07.022.869 I slot print_timing: id  2 | task 292 | n_decoded =   1528, tg =   3.46 t/s, tg_3s =   3.35 t/s
13.10.304.447 I slot print_timing: id  2 | task 292 | n_decoded =   1539, tg =   3.46 t/s, tg_3s =   3.35 t/s
13.13.331.062 I slot print_timing: id  2 | task 292 | n_decoded =   1549, tg =   3.46 t/s, tg_3s =   3.30 t/s
13.16.387.985 I slot print_timing: id  2 | task 292 | n_decoded =   1559, tg =   3.46 t/s, tg_3s =   3.27 t/s
13.19.686.410 I slot print_timing: id  2 | task 292 | n_decoded =   1570, tg =   3.45 t/s, tg_3s =   3.33 t/s
13.22.688.521 I slot print_timing: id  2 | task 292 | n_decoded =   1580, tg =   3.45 t/s, tg_3s =   3.33 t/s
13.25.691.264 I slot print_timing: id  2 | task 292 | n_decoded =   1590, tg =   3.45 t/s, tg_3s =   3.33 t/s
13.28.985.265 I slot print_timing: id  2 | task 292 | n_decoded =   1601, tg =   3.45 t/s, tg_3s =   3.34 t/s
13.31.987.531 I slot print_timing: id  2 | task 292 | n_decoded =   1611, tg =   3.45 t/s, tg_3s =   3.33 t/s
13.35.258.157 I slot print_timing: id  2 | task 292 | n_decoded =   1622, tg =   3.45 t/s, tg_3s =   3.36 t/s
13.38.293.844 I slot print_timing: id  2 | task 292 | n_decoded =   1632, tg =   3.45 t/s, tg_3s =   3.29 t/s
13.41.361.588 I slot print_timing: id  2 | task 292 | n_decoded =   1642, tg =   3.45 t/s, tg_3s =   3.26 t/s
13.44.366.202 I slot print_timing: id  2 | task 292 | n_decoded =   1652, tg =   3.45 t/s, tg_3s =   3.33 t/s
13.47.374.761 I slot print_timing: id  2 | task 292 | n_decoded =   1662, tg =   3.45 t/s, tg_3s =   3.32 t/s
13.50.393.027 I slot print_timing: id  2 | task 292 | n_decoded =   1672, tg =   3.45 t/s, tg_3s =   3.31 t/s
13.53.439.214 I slot print_timing: id  2 | task 292 | n_decoded =   1682, tg =   3.45 t/s, tg_3s =   3.28 t/s
13.56.516.501 I slot print_timing: id  2 | task 292 | n_decoded =   1692, tg =   3.44 t/s, tg_3s =   3.25 t/s
13.59.609.926 I slot print_timing: id  2 | task 292 | n_decoded =   1702, tg =   3.44 t/s, tg_3s =   3.23 t/s
13.59.921.983 I slot print_timing: id  2 | task 292 | prompt eval time =  140567.49 ms / 11995 tokens (   11.72 ms per token,    85.33 tokens per second)
13.59.921.994 I slot print_timing: id  2 | task 292 |        eval time =  494667.69 ms /  1703 tokens (  290.47 ms per token,     3.44 tokens per second)
13.59.921.997 I slot print_timing: id  2 | task 292 |       total time =  635235.18 ms / 13698 tokens
13.59.922.000 I slot print_timing: id  2 | task 292 |    graphs reused =          0
13.59.922.076 I slot      release: id  2 | task 292 | stop processing: n_tokens = 13697, truncated = 0
13.59.922.125 I srv  update_slots: all slots are idle

2.8 程序运行结果

python3 test-video.py 
模型回答:
<think>
用户现在需要详细描述这个视频内容,结合所有帧来分析场景和动作。首先,得先看每帧的画面元素、人物互动、字幕信息,然后整合起来。

首先看场景:整体是室内环境,看起来是家庭住宅,背景有厨房区域(能看到微波炉、橱柜、灯光),还有客厅部分(窗帘、置物架)。光线偏暗,色调偏冷,可能是夜晚或室内灯光较暗的场景,营造紧张氛围。

人物:一男一女,男性穿着灰色带深色拼接的Polo衫,女性穿深色上衣(棕色或深紫色)。两人面对面,距离很近,明显在激烈争吵。

然后看每帧的字幕和动作变化:

第一帧(最上面):男性侧对镜头,女性背对镜头,男性表情严肃,女性在看男性。背景有置物架(微波炉、杯子等)。

第二帧:男性双手抬起,似乎在解释或争辩,字幕“我也想跟你好好过 I tried to have a life with you,” 表明男性在表达想好好过日子的愿望,但可能被对方误解。

第三帧:女性开口说话,字幕“就你干的这些破事儿 After you got us into this shit,” 女性指责男性做了坏事导致问题,语气愤怒。

第四帧:女性继续指责,“你娶我就是看上我们家钱了 You married me for our money.” 男性反驳,“你没花我家钱吗 Haven't you?” 两人在金钱问题上争执,涉及婚姻动机和经济纠纷。

第五帧:女性说“你平常应酬送的礼 the fancy gifts for business networking,” 指责男性用家庭钱应酬送礼,男性回应“我赚钱的时候 When I made money,” 可能想解释自己赚钱时的情况。

第六帧:男性情绪激动,手指指向女性,字幕“你不要动手 Hands off!” 说明男性可能觉得女性要动手,或者自己在警告对方不要动手,动作上男性伸出手臂指向,姿态紧张。

第七帧:男性站在女性对面,女性侧身,男性表情严肃,可能在继续争执,背景有窗帘和置物架,灯光较暗。

现在整合这些信息:

场景是家庭室内,可能是客厅与厨房交界处,夜晚或昏暗环境,氛围紧张。两人是夫妻关系(从对话“你娶我”等判断),正在激烈争吵,核心围绕婚姻动机(是否为钱结婚)、经济纠纷(应酬送礼、花钱问题)。

动作细节:男性多次用手势强调(抬手、指向),表情从解释到愤怒;女性表情激动,说话时语气强硬,指责对方。字幕显示对话内容逐步升级,从男性表达愿望,到女性指责行为,再到金钱动机、应酬花费,最后男性警告“不要动手”,显示冲突升级,可能有肢体冲突风险。

需要确认每帧的连续性:对话顺序是男性先说想好好过,女性指责他搞砸了,说他娶自己为钱,男性反问她没花家里的钱吗,女性说他应酬送礼,男性说赚钱时的情况,然后男性警告不要动手。整个过程是争吵的递进,情绪从争执到紧张升级。

还要注意环境细节:背景有微波炉、置物架、窗帘,厨房区域有灯光,说明是居家环境,可能是晚餐后或夜间在家发生争执。

总结时要涵盖场景、人物关系、对话内容、动作变化、情绪发展。
</think>

这段视频呈现了一对夫妻在**居家室内环境**中爆发激烈争吵的场景,通过多帧画面的连续互动,展现了冲突的递进过程与情绪张力。以下结合场景、人物动作、对话字幕等细节进行详细解析:  


### **一、场景与环境**  
画面背景为典型的**家庭住宅空间**,左侧可见带有圆形花纹的窗帘、木质置物架(架上摆放微波炉、纸盒、杯具等生活用品),右侧是厨房区域(能看到橱柜、水槽、灯光照明的台面),整体光线偏暗、色调偏冷(蓝绿色调为主),营造出**夜晚或室内昏暗时段**的压抑氛围,暗示冲突发生在私密且情绪敏感的居家场景中。  


### **二、人物关系与核心冲突**  
画面中一男一女面对面近距离对峙,从对话字幕“你娶我就是看上我们家钱了”“I tried to have a life with you”等信息可判断,两人是**夫妻关系**,争吵核心围绕**婚姻动机、经济纠纷**展开,情绪从争执逐渐升级为激烈对抗。  


### **三、动作与对话的递进过程**  
通过多帧画面的连续性,可梳理出争吵的逻辑链条:  

1. **男性试图解释(第一、二帧)**  
   - 男性身穿灰色拼接Polo衫,初始姿态略显无奈,随后**双手抬起、手掌张开**,配合字幕“我也想跟你好好过 / I tried to have a life with you,” 表达自己曾希望婚姻平稳的意愿,动作带有“辩解”“恳求”的意味。  

2. **女性激烈指责(第三、四帧)**  
   - 女性留黑色长发,身穿深色上衣,面对男性时**面部紧绷、语气尖锐**。  
   - 第三帧字幕“就你干的这些破事儿 / After you got us into this shit,” 直接将矛盾归咎于男性行为;  
   - 第四帧进一步升级指责:“你娶我就是看上我们家钱了 / You married me for our money.” 男性随即反问“你没花我家钱吗 / Haven't you?”,双方陷入**金钱动机的互相攻讦**——女性认为婚姻是“为钱”,男性则质疑妻子未承担家庭经济责任。  

3. **经济纠纷细化(第五帧)**  
   - 女性继续列举具体矛盾:“你平常应酬送的礼 / the fancy gifts for business networking,” 指责男性用家庭资金进行社交应酬;  
   - 男性回应“我赚钱的时候 / When I made money,” 暗示自己曾为家庭经济付出,试图扭转“被指责为贪财”的叙事。  

4. **冲突升级至肢体警告(第六、七帧)**  
   - 男性情绪爆发,**身体前倾、手指直指女性**,字幕“你不要动手 / Hands off!” 显示他既可能是在警告对方“不要动手”,也可能因愤怒而做出防御性动作;  
   - 第七帧中男性姿态紧绷、面部表情严肃,女性侧身回避但未退让,双方距离极近却充满敌意,暗示冲突已逼近肢体对抗的临界点。  


### **四、情绪与氛围的传递**  
- **男性**:从“试图沟通”到“情绪失控”,动作从“抬手解释”变为“手指指责”,体现从理性辩解到愤怒防御的心理转变;  
- **女性**:全程以**尖锐语气、直视对方**的姿态主导指责,表情中透露出失望与愤怒,将矛盾聚焦于“婚姻动机”与“经济不公”;  
- **环境烘托**:昏暗光线、居家私密空间与厨房的冷光形成对比,既强化了“家庭矛盾”的真实感,又通过冷色调放大了情绪的冰冷与紧张。  


### **总结**  
这段视频通过**近距离对峙、递进式对话、肢体语言变化**,生动刻画了一对夫妻因**婚姻信任崩塌、经济纠纷**引发的激烈争吵。场景细节(居家环境、生活用品)与情绪张力(从争执到警告)结合,精准传递出亲密关系中因金钱、动机等核心问题爆发的冲突本质,具有强烈的现实感与戏剧张力。

三、踩坑及解决方案

3.1 cooperative_groups头文件编译失败

错误现象

/opt/dtk/include/hip/amd_detail/cooperative_groups/details/sync.h:70:62: 
error: invalid input constraint 'l' in asm

原因分析:DTK 26.04的cooperative_groups头文件中使用了内联汇编约束'l',但gfx928架构不支持该约束。

解决方案:修改/opt/dtk-26.04/include/hip/amd_detail/cooperative_groups/details/info.h,将_CG_ASM_PTR_CONSTRAINT"l"改为"r"

3.2 __any重复定义错误

错误现象

error: redefinition of '__any'

原因分析__any是HIP内建函数,而cooperative_groups头文件中的__any_sync尝试调用__any,导致冲突。

解决方案:在info.h中注释掉__any_sync函数。

3.3 Launch params超限错误

错误现象

Launch params (1024, 1, 1) are larger than launch bounds (256) for kernel rms_norm_f32

原因分析:gfx928硬限制线程块最大为256,但llama.cpp的GPU内核硬编码为1024线程。

解决方案:将所有内核中的1024线程块改为256,同时修改对应的静态断言和条件判断。

3.4 mul_mat_f内核无设备代码

错误现象

ERROR: HIP kernel mul_mat_f has no device code compatible with HIP arch 1300

原因分析mul_mat_f依赖于MFMA(Matrix FMA)指令,但gfx928上该指令未正确启用或编译。

解决方案:在mmf.cu中强制ggml_cuda_should_use_mmf返回false,并在编译时添加-DGGML_CUDA_FORCE_MMQ=ON,强制使用MMQ矩阵乘法内核。

3.5 Flash Attention内核缺失

错误现象

ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300

原因分析:Flash Attention的MFMA版本同样缺失gfx928的设备代码。

解决方案:启动服务时添加--flash-attn off参数,禁用Flash Attention,使用普通的注意力计算。

3.6 性能问题

现象:关闭Flash Attention后,BF16模型的推理速度约为4.5 tokens/s。

建议

(1)使用Q4_K_M量化模型可显著提升速度并降低显存占用
(2)后续可尝试在fattn-mma-f16.cuh中手动启用AMD_MFMA_AVAILABLE并添加编译选项(目前海光DTK环境暂不支持)
(3)考虑切换到已适配的atomic-llama-cpp-turboquant分支(目前暂不支持定制分支)

四、总结

本文详细记录了在海光K100_AI(DCU,gfx928架构)上编译快手定制版llama.cpp并部署Keye-VL-2.0-30B-A3B多模态大模型的完整过程。

核心经验总结

4.1 llama.cpp是海光DCU上部署新模型的现实选择:在vLLM、SGLang、Transformers等主流框架对海光DCU和新模型的支持滞后的情况下,llama.cpp凭借其轻量级和跨平台特性,成为了唯一可行的部署路径。

4.2 国产GPU适配的核心挑战在于硬件限制与软件假设的冲突:gfx928的256线程硬限制与llama.cpp的1024线程默认配置之间的矛盾,需要通过大量源码修改来解决。

4.3 系统级头文件修改是绕不开的坎:DTK的cooperative_groups头文件与gfx928的兼容性问题,需要通过修改系统头文件来绕过。

4.4 MMQ是救急良药:当MFMA等高级指令无法使用时,强制使用MMQ矩阵乘法内核可以保证程序稳定运行。

4.5 性能与兼容性的权衡:关闭Flash Attention和禁用MFMA虽然影响了性能,但换来了稳定运行。后续可通过量化模型、启用MTP投机解码等方式优化速度。

展望:随着国产GPU生态的不断完善,期待海光DCU对vLLM、SGLang等主流框架的支持能更加完善,未来能够以更少的适配工作完成大模型的部署。

本文所有操作均在Ubuntu 22.04 + DTK 26.04 + 海光K100_AI环境下验证通过。如有问题,欢迎在评论区交流讨论。

更多推荐