CANN/asc-devkit:__cvta_local_to_generic地址空间转换函数

📅 2026/7/17 8:34:58 👁️ 阅读次数 📝 编程学习
CANN/asc-devkit:__cvta_local_to_generic地址空间转换函数

__cvta_local_to_generic

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

功能说明

将栈空间的地址值转换为对应的指针并返回。

函数原型

void* __cvta_local_to_generic(size_t rawbits)

参数说明

表1参数说明

参数名输入/输出描述
rawbits输入源操作数。

返回值说明

指向输入栈空间的指针。该接口不校验输入是否为有效栈空间地址。异常场景示例说明如下:

输入场景返回值
rawbits0返回指针使用__isLocal判断为1,但该指针不一定是可安全访问的栈空间地址。
rawbits为全1返回指针使用__isUbuf判断为1,说明非法rawbits的分类结果不可依赖。

约束说明

SIMD与SIMT混合编程场景不支持使用该接口。

需要包含的头文件

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

#include "simt_api/device_functions.h"

调用示例

SIMT编程场景:

__global__ __launch_bounds__(1024) void kernel___cvta_local_to_generic(uint32_t* dst, uint32_t* src) { int idx = threadIdx.x + blockIdx.x * blockDim.x; uint32_t data_ptr[]={10,11}; size_t addr = __cvta_generic_to_local(data_ptr); void* ptr = __cvta_local_to_generic(addr); dst[idx] = *(uint32_t*)ptr; }

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

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考