企业官网建设流程全解析

旅游网站建设前的市场分析,wordpress特定账户注册,安徽省建设工程安全协会网站,wordpress统计工具训练营简介 2025年昇腾CANN训练营第二季#xff0c;基于CANN开源开放全场景#xff0c;推出0基础入门系列、码力全开特辑、开发者案例等专题课程#xff0c;助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证#xff0c;即可领取精美证书#xff0c;完成…训练营简介 2025年昇腾CANN训练营第二季基于CANN开源开放全场景推出0基础入门系列、码力全开特辑、开发者案例等专题课程助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证即可领取精美证书完成社区任务更有机会赢取华为手机平板、开发板等大奖。报名链接https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro摘要在算子开发初期开发者往往习惯将 Shape 写死在代码中Hardcode但这在动态输入的真实推理场景下寸步难行。CANN 架构采用了独特的Host Tiling Kernel Computing分离模式。本文将深入解析这一设计哲学教你如何利用 Host 侧的 CPU 算力“运筹帷幄”计算出最优切分策略指挥 Device 侧的 NPU 灵活应对任意尺寸的输入数据。前言为什么 Kernel 不能自己决定怎么切在 CUDA 编程中我们习惯在核函数Kernel内部通过blockIdx和threadIdx动态计算数据的偏移量。但在昇腾 Ascend C 开发中你会发现一个奇怪的现象算子的切分逻辑Tiling必须在 Host 侧写而不能全塞进 Kernel 里。这是由达芬奇架构Da Vinci Architecture的特性决定的AI Core是极致的计算怪兽它的 Scalar 单元虽然能做逻辑控制但极其宝贵应该专注于流水线编排。Host CPU擅长复杂的逻辑运算和分支判断。因此CANN 的设计哲学是让 CPU 负责“分蛋糕”计算 Tiling 参数让 NPU 负责“吃蛋糕”执行计算。只有掌握了动态 Tiling你的算子才算真正具备了通用性。一、 核心图解TilingData 的“锦囊妙计”动态 Shape 的核心在于TilingData结构体。它就像是 Host 军师交给 Device 将军的一个“锦囊”。Host 侧根据输入的实际 Shape例如[32, 1024]或[64, 2048]和硬件资源UB 大小、Core 数量计算出本次运行需要的核心数BlockDim、每个核心处理的数据量TileLength等参数。传输将这些参数打包成TilingData结构体拷贝到 Device 侧的 Workspace。Device 侧Kernel 启动时打开“锦囊”读取参数直接按指令干活不再进行复杂的切分计算。二、 关键策略如何优雅地“切蛋糕”一个优秀的 Tiling 策略需要同时考虑多核负载均衡和UB 内存限制。2.1 策略一多核切分 (Block Dim Strategy)假设输入数据总量为TotalLength可用核心数为aicore_num。平均分配BlockLength TotalLength / aicore_num。尾数处理如果除不尽前remainder个核心多干一点后面的核心少干一点。CANN 最佳实践不要让任何一个核心闲着Idle。算力空转是最大的犯罪。2.2 策略二UB 切分 (L1/UB Strategy)单个核心拿到的数据BlockLength可能依然很大例如 1MB而 UB 只有 256KB。这就需要进行二次切分。双缓冲约束为了开启 Ping-Pong 流水线我们需要将 UB 划分为两块每块大小为UB_Size / 2。对齐约束这是最容易踩的坑昇腾硬件要求数据搬运地址必须32 字节对齐32 Bytes Alignment。如果切分出的TileLength不是 32B 的倍数例如 float16 下不是 16 的倍数DataCopy 就会报错或读写越界。三、 实战手写动态 Tiling 逻辑3.1 定义 TilingData (Host/Device 共用头文件)这是双方的通信协议。// common.h struct TilingData { uint32_t totalLength; // 总数据量 uint32_t tileNum; // 每个核内切分多少次 uint32_t tileLength; // 每次切分的数据长度 uint32_t lastTileLength; // 最后一个切分的长度处理尾块 uint32_t coreDataNum; // 每个核处理的总数据量 };3.2 Host 侧计算逻辑 (Tiling 函数)// host_tiling.cpp void ComputeTiling(const uint32_t totalLength, TilingData tiling) { // 1. 获取硬件信息 auto ascendInfo platform_ascend::GetPlatformInfo(); uint32_t coreNum ascendInfo.GetCoreNum(); uint32_t ubSize ascendInfo.GetUbSizeInBytes(); // 2. 计算每个核分多少 (简化版假设能整除) tiling.coreDataNum totalLength / coreNum; // 3. 计算核内分块 // 预留一部分 UB 给系统假设可用 80% uint32_t maxTileBytes ubSize * 0.8 / 2; // 双缓冲除以2 uint32_t maxDataNum maxTileBytes / sizeof(half); // 必须满足 32B 对齐 (16个 half) // ALIGN_DOWN 是向下对齐宏 uint32_t alignedDataNum ALIGN_DOWN(maxDataNum, 16); tiling.tileLength alignedDataNum; tiling.tileNum tiling.coreDataNum / tiling.tileLength; tiling.lastTileLength tiling.coreDataNum % tiling.tileLength; // ... 处理 lastTileLength 为 0 的情况 ... }3.3 Kernel 侧执行逻辑// kernel.cpp extern C __global__ __aicore__ void MyAdd(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { // 1. 获取 TilingData // GET_TILING_DATA 是 CANN 提供的宏自动从 Global Memory 读取参数解析到结构体 GET_TILING_DATA(tilingData, tiling); // 2. 根据 Host 算好的参数干活 // 不需要再思考 我该分多少直接用 tilingData.tileLength KernelAdd op; op.Init(x, y, z, tilingData.totalLength, tilingData.tileLength, ...); op.Process(); }四、 进阶非对齐 Shape 的终极方案如果输入 Shape 极其刁钻比如Length 31怎么都凑不齐 32B 对齐怎么办CANN 提供了UB Padding技巧搬运时使用DataCopy的pad模式或者手动搬运多一点数据凑齐 32B。计算时Mask 掉无效数据或者多算一点只要不写回越界。写回时利用 DataCopy 的高级参数只写回有效的 31 个数据。虽然这会增加逻辑复杂度但为了通用性这是必经之路。五、 总结动态 Tiling 是 Ascend C 算子从 Demo 走向 Product 的分水岭。分离思想CPU 负责策略NPU 负责执行。对齐意识时刻谨记 32 Bytes 对齐这是达芬奇架构的底线。极致利用通过精细计算 UB 占用尽可能把 UB 塞满减少搬运次数。当你习惯了用 TilingData 传递参数而不是在 Kernel 里写#define LENGTH 1024时你就掌握了 CANN 开发的精髓。本文基于昇腾 CANN 8.0 架构特性编写。