Ascend C SIMT Cache一致性接口文档

asc_dcci_single

【免费下载链接】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 训练系列产品：不支持

功能说明

该接口用于刷新指定地址所在的Cache Line，保证数据读取时Cache的一致性。

在SIMT编程中，写入数据时会立即写入Global Memory，使其他核可见，因此不存在核间一致性问题。当从Global Memory读取数据时，该数据可能已被其他核修改，此时应使用dcci接口直接访问Global Memory，以获取最新数据。

函数原型

inline void asc_dcci_single(void *dst)

参数说明

表1参数说明

返回值说明

无

约束说明

输入必须为有效的Global Memory地址。传入非Global Memory地址或无效地址时，不能保证Cache刷新行为有效。

需要包含的头文件

使用该接口需要包含"simt_api/device_functions.h"头文件。

#include "simt_api/device_functions.h"

调用示例

• SIMT编程场景：

  __global__ __launch_bounds__(1024) void kernelDcci_single(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_single((void*)(data)); // 从Global Memory读标志位 } output[0] = data[1]; // 数据就绪，获取数据 } }

• SIMD与SIMT混合编程场景：

  __simt_vf__ __launch_bounds__(1024) inline void kernelDcci_single(__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_single((__gm__ void*)(data)); // 从Global Memory读标志位 } output[0] = data[1]; // 数据就绪，获取数据 } }

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit