CANN/asc-devkit SIMT访存函数asc_dcci_entire
asc_dcci_entire【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品不支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明该接口用于刷新核内的整个Data Cache保证数据读取时Cache的一致性。在SIMT编程中写入数据时会立即写入Global Memory使其他核可见因此不存在核间一致性问题。当从Global Memory读取数据时该数据可能已被其他核修改此时应使用dcci接口直接访问Global Memory以获取最新数据。请注意该接口耗时较大性能敏感的场景慎用。函数原型inline void asc_dcci_entire(void *dst)参数说明表1参数说明参数名输入/输出描述dst输入Global Memory的地址。返回值说明无约束说明输入必须为有效的Global Memory地址。传入非Global Memory地址或无效地址时不能保证Cache刷新行为有效。需要包含的头文件使用该接口需要包含simt_api/device_functions.h头文件。#include simt_api/device_functions.h调用示例SIMT编程场景__global__ __launch_bounds__(1024) void kernelDcci_entire(uint32_t* data, uint32_t* output) { // 线程块0的线程0进行数据写操作 if (blockIdx.x 0 threadIdx.x 0) { data[1] 10; // 写入数据 asc_threadfence(); // 确保数据写入完成 data[0] 1; // 标志位数据已就绪 } // 线程块1的线程0进行数据读操作 if (blockIdx.x 1 threadIdx.x 0) { while (data[0] ! 1) { asc_dcci_entire((void*)(data)); // 从Global Memory读标志位 } output[0] data[1]; // 数据就绪获取数据 } }SIMD与SIMT混合编程场景__simt_vf__ __launch_bounds__(1024) inline void kernelDcci_entire(__gm__ uint32_t* data, __gm__ uint32_t* output) { // 线程块0的线程0进行数据写操作 if (blockIdx.x 0 threadIdx.x 0) { data[1] 10; // 写入数据 asc_threadfence(); // 确保数据写入完成 data[0] 1; // 标志位数据已就绪 } // 线程块1的线程0进行数据读操作 if (blockIdx.x 1 threadIdx.x 0) { while (data[0] ! 1) { asc_dcci_entire((__gm__ void*)(data)); // 从Global Memory读标志位 } output[0] data[1]; // 数据就绪获取数据 } }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
