English | 中文版

附录 C：300 个 MultiKernelBench 内核的漏洞分析

MultiKernelBench 的 300 个内核涵盖 15 个类别。如果按照标准 AscendC C++ 方式实现，每个内核都会继承 GM_ADDR/LocalTensor/FreeTensor API 的结构性漏洞模式。我们系统分类哪些模式影响哪些内核类别，统计暴露面，并展示最高风险的 C++ 与 ascend-rs 对比。

C.1 漏洞模式分布

关键发现：每个 AscendC C++ 内核在结构上都暴露于 V1（类型擦除）、V4（释放后使用）、V5（双重释放）和 V6（缺失同步），因为这些是 API 本身的属性，而非特定算法的问题。算法性漏洞（V2、V3）影响的子集取决于内核是否使用逐元素索引访问或多核偏移算术。

C.2 最高风险类别：索引操作（12 个内核）

索引内核（gather、scatter、scatter_add、index_select、index_copy、index_add、embedding、masked_fill、inplace_update、take_along_dim、argmax、argmin）是最高风险类别，因为它们同时组合了全部六种漏洞模式：

• V1：GM_ADDR 擦除张量元素类型
• V2：用户提供的索引值无边界检查地访问任意偏移
• V3：idx * row_len + j 对大张量可能溢出 uint32_t
• V4/V5：分块实现使用 FreeTensor 生命周期管理
• V6：需要 DMA 与计算之间的同步

C++ AscendC gather（存在漏洞）：

#include "kernel_operator.h"

// GM_ADDR 擦除所有类型信息——调用者可以传入任何数据类型
extern "C" __global__ __aicore__
void gather(GM_ADDR input, GM_ADDR index, GM_ADDR output, GM_ADDR len_buf) {
    uint32_t n = *((__gm__ uint32_t *)len_buf);
    // V1：从 GM_ADDR 手动转换——无编译期类型安全
    __gm__ float *in_ptr = (__gm__ float *)input;
    __gm__ uint32_t *idx_ptr = (__gm__ uint32_t *)index;
    __gm__ float *out_ptr = (__gm__ float *)output;

    for (uint32_t i = 0; i < n; i++) {
        uint32_t idx = idx_ptr[i];
        // V2：idx 无边界检查——攻击者控制的索引
        // 可读取 GM 地址空间内的任意内存
        out_ptr[i] = in_ptr[idx];  // 若 idx >= input_len 则越界
    }
}

ascend-rs gather（已缓解）：

#![allow(unused)]
fn main() {
#[ascend_std::aiv_kernel]
pub unsafe fn gather(
    input: *const f32,   // V1 已缓解：类型化指针，非 GM_ADDR
    index: *const u32,
    output: *mut f32,
    len: *const u32,
) {
    unsafe {
        let n = *len;
        let mut i = 0u32;
        loop {
            if i >= n { break; }  // 循环边界显式表达
            let idx = *index.wrapping_add(i as usize);
            // V2：wrapping_add 显式表达指针算术语义
            // V3：无整数溢出——每个偏移独立转换
            *output.wrapping_add(i as usize) = *input.wrapping_add(idx as usize);
            i = i + 1;
        }
        // V4/V5：无 FreeTensor——缓冲区 ID 自动管理
        // V6：无 DMA/计算分离——标量操作直接访问 GM
    }
}
}

C.3 高风险类别：卷积内核（34 个内核）

卷积内核具有深层嵌套循环和复杂的多维索引算术（oc * in_ch * k_h * k_w + ic * k_h * k_w + kh * k_w + kw）。索引表达式中的单个维度错误会静默读取错误内存。

C++ AscendC conv2d 索引计算（存在漏洞）：

// V2+V3：6层嵌套索引算术——极易弄错某个维度
for (int oc = 0; oc < out_ch; oc++) {
    for (int oh = 0; oh < out_h; oh++) {
        for (int ow = 0; ow < out_w; ow++) {
            float sum = 0.0f;
            for (int ic = 0; ic < in_ch; ic++) {
                for (int kh = 0; kh < k_h; kh++) {
                    for (int kw = 0; kw < k_w; kw++) {
                        int ih = oh * stride + kh * dilation;
                        int iw = ow * stride + kw * dilation;
                        // V3：32位乘法链可能溢出
                        int in_idx = ic * in_h * in_w + ih * in_w + iw;
                        int w_idx = oc * in_ch * k_h * k_w
                                  + ic * k_h * k_w + kh * k_w + kw;
                        // V2：无边界检查——若 ih >= in_h 或 iw >= in_w，
                        // 则从 GM 越界读取
                        sum += (float)inLocal.GetValue(in_idx)
                             * (float)wLocal.GetValue(w_idx);
                    }
                }
            }
            outLocal.SetValue(oc * out_h * out_w + oh * out_w + ow, sum);
        }
    }
}

ascend-rs conv2d（已缓解）：

#![allow(unused)]
fn main() {
#[ascend_std::aiv_kernel]
pub unsafe fn conv_standard_2d(
    input: *const f32, weight: *const f32, output: *mut f32,
    params: *const u32,  // [in_ch, out_ch, in_h, in_w, k_h, k_w, stride, dilation]
) {
    unsafe {
        // 所有参数从类型化指针读取——无 GM_ADDR 转换
        let in_ch = *params;
        let out_ch = *params.wrapping_add(1);
        // ...（读取其余参数）
        let out_h = (in_h - (k_h - 1) * dilation - 1) / stride + 1;
        let out_w = (in_w - (k_w - 1) * dilation - 1) / stride + 1;

        let mut oc = 0u32;
        loop {
            if oc >= out_ch { break; }
            // ...显式边界的嵌套循环...
            let ih = oh * stride + kh * dilation;
            let iw = ow * stride + kw * dilation;
            // V3 已缓解：通过 `as usize` 显式表达 wrapping 语义
            // 调试构建溢出时 panic，发布构建有意 wrapping
            let in_idx = (ic * in_h * in_w + ih * in_w + iw) as usize;
            let w_idx = (oc * in_ch * k_h * k_w
                       + ic * k_h * k_w + kh * k_w + kw) as usize;
            sum = sum + *input.wrapping_add(in_idx) * *weight.wrapping_add(w_idx);
            // V4/V5：无需 FreeTensor
            // V6：无 DMA——标量 GM 访问
        }
    }
}
}

C.4 高风险类别：融合算子（100 个内核）

融合内核（matmul+activation、conv+norm+activation 等）串联多个流水线阶段。在 C++ 中，每个阶段都需要各自的 AllocTensor/FreeTensor/pipe_barrier——遗漏任何一个都会产生静默数据损坏。

C++ 融合 matmul+sigmoid（存在漏洞）：

// 融合 matmul + sigmoid：C = sigmoid(A * B)
// V4：分配/释放 4 个张量——每个都是释放后使用的机会
// V5：融合变体之间的复制粘贴可能重复 FreeTensor
// V6：3 次流水线转换（DMA->cube, cube->vector, vector->DMA）
//     ——每次都需要 pipe_barrier，遗漏任何一个 = 读取过期数据

AscendC::LocalTensor<half> aLocal = inQueueA.AllocTensor<half>();
AscendC::DataCopy(aLocal, aGm, m * k);
inQueueA.EnQue(aLocal);
// V6：此处需要 DMA -> cube 的屏障
aLocal = inQueueA.DeQue<half>();

// ...矩阵乘法...

inQueueA.FreeTensor(aLocal);
// V4：aLocal 句柄仍然有效——意外读取能编译和运行

AscendC::LocalTensor<float> cLocal = outQueue.AllocTensor<float>();
// V6：此处需要 cube -> vector 的屏障
AscendC::Muls(cLocal, cLocal, -1.0f, total);  // sigmoid 步骤 1
AscendC::Exp(cLocal, cLocal, total);            // sigmoid 步骤 2
// V6：310P 上同缓冲区就地链式操作需要操作间屏障
AscendC::Adds(cLocal, cLocal, 1.0f, total);    // sigmoid 步骤 3
AscendC::Reciprocal(cLocal, cLocal, total);     // sigmoid 步骤 4
outQueue.FreeTensor(cLocal);

ascend-rs 融合 matmul+sigmoid（已缓解）：

#![allow(unused)]
fn main() {
#[ascend_std::aiv_kernel]
pub unsafe fn fused_matmul_sigmoid(
    a: *const u16, b: *const u16, c: *mut f32, dims: *const u32,
) {
    unsafe {
        let m = *dims;
        let k = *dims.wrapping_add(1);
        let n = *dims.wrapping_add(2);

        // V6 已缓解：matmul_f16 内部处理 DMA+cube
        ascend_std::kernel_ops::matmul_f16(c, a, b, m, k, n);
        ascend_std::ascend_pipe_barrier();  // 显式、可见

        let total = m * n;
        let buf_c = ascend_std::ascend_buf_alloc(total);
        ascend_std::ascend_buf_load_f32(buf_c, c as *const f32, total);
        ascend_std::ascend_pipe_barrier();  // 显式、可见

        // V6 已缓解：sigmoid_f32 包含所有内部屏障
        // (muls -> barrier -> exp -> barrier -> adds -> barrier -> reciprocal)
        ascend_std::kernel_ops::sigmoid_f32(buf_c, buf_c, total);

        ascend_std::ascend_pipe_barrier();  // 显式、可见
        ascend_std::ascend_buf_store_f32(c, buf_c, total);
        // V4/V5：无 FreeTensor——buf_c 自动管理
    }
}
}

C.5 漏洞统计：300 个内核 x 6 种模式

C.6 ascend-rs 如何消除每种模式

净效果：在 300 个内核总共 1,633 个漏洞暴露中，ascend-rs 在 API/类型层面消除了 1,500 个（V1、V4、V5 完全消除；V6 从每操作减少到每内核）。剩余的 133 个越界暴露（V2）通过将逐元素访问替换为整向量操作来缓解，但标量回退内核中的 unsafe 指针算术仍需程序员负责。