AtomicCas【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品xAtlas A2 训练系列产品 / Atlas A2 推理系列产品x功能说明调用该接口后可在指定GM地址上进行原子比较如果和value1相等则把value2的值赋值到GM上如果和value1不相等则GM上的值不变。函数原型template typename T __aicore__ inline T AtomicCas(__gm__ T *address, T value1, T value2)参数说明表 1模板参数说明参数名描述T操作数数据类型。Ascend 950PR/Ascend 950DT支持的数据类型为uint32_t/uint64_t表 2参数说明参数名输入/输出描述address输入输入GM的地址。value1/value2输入标量值支持数据类型和address指向的数据类型保持一致。返回值说明GM地址上做原子操作前的数据。约束说明原子操作涉及标量计算如果标量计算单元和搬运单元MTE2/MTE3在读写GM时存在数据依赖开发者需要根据实际情况插入同步。调用示例extern C __global__ __aicore__ void atomic_simple_kernel(__gm__ uint8_t* dst uint32_t dataSize) { // ... dst_global.SetGlobalBuffer(reinterpret_cast__gm__ T *(dst_gm), dataSize); LocalTensorT dstLocal inQueueX.AllocTensorT(); uint32_t value1 1; uint32_t value2 2; uint32_t a AscendC::AtomicCas(reinterpret_cast__gm__ int32_t *(dst), value1, value2); // 先执行完原子操作之后才能进行搬运操作有数据依赖需要手动插入MTE2等待Scalar的同步 event_t eventIdSToMte2 static_castevent_t(GetTPipePtr()-AllocEventIDHardEvent::S_MTE2()); SetFlagHardEvent::S_MTE2(eventIdSToMte2); WaitFlagHardEvent::S_MTE2(eventIdSToMte2); DataCopy(dstLocal, dst_global, dataSize); // ... }假设上述函数在3个核上执行核1、核2、核3依次调度结果示例如下原GM数据dst: [1,1,1,1,1,...,1] 核1 原子计算后GM数据dst: [2,1,1,1,1,...,1] 返回值 a: 1 核2 原子计算后GM数据dst: [2,1,1,1,1,...,1] 返回值 a: 2 核3 原子计算后GM数据dst: [2,1,1,1,1,...,1] 返回值 a: 2【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考