海光K100_AI环境编译快手定制版llama.cpp的艰难探索：部署测试Keye-VL-2.0-30B-A3B多模态大模型（含完整过程及代码）

一、引言

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 环境说明

本文的部署环境如下：

二、部署测试详细步骤

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
7template <int block_size>
8static __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
44template <int block_size>
45static __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
79template <int block_size, bool do_multiply = false, bool do_add = false>
80__launch_bounds__(256, 1)
81static __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
161template <int block_size>
162static __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
248template <int block_size>
249static __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
284static 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
297static 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
308static 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
330static 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
414static 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
424static 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
439void 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
463void 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
482void 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
506void 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
566void 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
654void 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
680void 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环境下验证通过。如有问题，欢迎在评论区交流讨论。