4. 向量加法内核实现

任务目标

在寒武纪MLU上实现一个向量加法内核：输入两个长度为N的浮点数组A、B，计算C = A + B。内核在MLU上并行执行，并将结果与CPU逐元素计算结果对比，验证正确性。

环境与依赖

• 寒武纪MLU开发环境（含cncc编译器、cnrt运行时）
• 已设置NEUWARE_HOME环境变量
• 单卡MLU设备（代码中默认使用设备0）

完整实现代码

// vector_add.cpp
#include <cnrt.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>

#define N 1024
#define THREADS 256

// MLU端内核：向量加法
__global__ void vectorAddKernel(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    // 1. 初始化MLU设备
    cnrtInit(0);
    cnrtDev_t dev;
    cnrtGetDeviceHandle(&dev, 0);
    cnrtSetCurrentDevice(dev);

    // 2. 分配主机内存并初始化数据
    float *h_a = (float*)malloc(N * sizeof(float));
    float *h_b = (float*)malloc(N * sizeof(float));
    float *h_c = (float*)malloc(N * sizeof(float));
    float *h_ref = (float*)malloc(N * sizeof(float));

    for (int i = 0; i < N; i++) {
        h_a[i] = (float)(rand() % 100) / 10.0f;
        h_b[i] = (float)(rand() % 100) / 10.0f;
    }

    // 3. 分配MLU设备内存
    float *d_a, *d_b, *d_c;
    cnrtMalloc(&d_a, N * sizeof(float));
    cnrtMalloc(&d_b, N * sizeof(float));
    cnrtMalloc(&d_c, N * sizeof(float));

    // 4. 主机数据拷贝到设备
    cnrtMemcpy(d_a, h_a, N * sizeof(float), CNRT_MEM_TRANS_DIR_HOST2DEV);
    cnrtMemcpy(d_b, h_b, N * sizeof(float), CNRT_MEM_TRANS_DIR_HOST2DEV);

    // 5. 配置内核启动参数并执行
    cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_BLOCK;
    cnrtDim3_t dim;
    dim.x = (N + THREADS - 1) / THREADS;
    dim.y = 1;
    dim.z = 1;
    cnrtKernelParamsBuffer_t params;
    cnrtKernelParamsBufferInit(&params);
    cnrtKernelParamsBufferAddParam(&params, &d_a, sizeof(d_a));
    cnrtKernelParamsBufferAddParam(&params, &d_b, sizeof(d_b));
    cnrtKernelParamsBufferAddParam(&params, &d_c, sizeof(d_c));
    cnrtKernelParamsBufferAddParam(&params, &N, sizeof(N));

    cnrtQueue_t queue;
    cnrtCreateQueue(&queue);
    cnrtLaunchKernel(vectorAddKernel, dim, &params, ktype, queue);
    cnrtSyncQueue(queue);

    // 6. 结果拷贝回主机
    cnrtMemcpy(h_c, d_c, N * sizeof(float), CNRT_MEM_TRANS_DIR_DEV2HOST);

    // 7. CPU参考计算
    for (int i = 0; i < N; i++) {
        h_ref[i] = h_a[i] + h_b[i];
    }

    // 8. 验证正确性
    float max_err = 0.0f;
    for (int i = 0; i < N; i++) {
        float err = fabsf(h_c[i] - h_ref[i]);
        if (err > max_err) max_err = err;
    }
    printf("Max error: %f\n", max_err);
    if (max_err < 1e-5) {
        printf("PASSED: MLU result matches CPU reference.\n");
    } else {
        printf("FAILED: MLU result deviates from CPU reference.\n");
    }

    // 9. 清理资源
    cnrtKernelParamsBufferDestroy(&params);
    cnrtDestroyQueue(queue);
    cnrtFree(d_a); cnrtFree(d_b); cnrtFree(d_c);
    free(h_a); free(h_b); free(h_c); free(h_ref);
    cnrtDestroy();
    return 0;
}

编译与运行

# 编译（生成MLU可执行文件）
cncc -o vector_add vector_add.cpp -I$NEUWARE_HOME/include -L$NEUWARE_HOME/lib64 -lcnrt

# 运行
./vector_add

验证方法与判定标准

• 程序输出Max error值，表示MLU计算结果与CPU参考结果的最大误差。
• 若输出包含PASSED且最大误差小于1e-5，则判定内核实现正确。
• 若输出FAILED，检查设备内存分配是否成功、队列同步是否完成、数据拷贝方向是否正确。

关键说明

• 内核使用blockIdx.x * blockDim.x + threadIdx.x计算全局索引，通过边界判断idx < n保证线程数不是N整数倍时仍安全。
• 本例使用CNRT_FUNC_TYPE_BLOCK类型，即每个线程独立执行一个元素加法，是MLU上最基础的并行模式。
• 随机数使用rand()生成，每次运行数据不同，但CPU与MLU使用相同输入，因此对比结果始终有效。