CANN/asc-devkit __isLocal函数文档

__isLocal产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品不支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明判断输入的指针是否指向栈空间的地址。函数原型unsigned int __isLocal(const void* ptr)参数说明表1参数说明参数名输入/输出描述ptr输入源操作数。返回值说明如果输入的指针指向栈空间的地址则返回1否则返回0。该接口根据输入指针的地址空间信息进行分类判断不校验该指针是否为可安全访问的有效地址。__isLocal返回1仅表示该指针被分类为栈空间地址不代表该地址一定可以安全访问。特殊场景说明如下| 输入场景 | 返回值 | | --- | --- | |ptr为有效Global Memory指针 | 0 | |ptr为有效Unified Buffer指针 | 0 | |ptr为nullptr| 0 | |ptr为伪造的低位地址如(void*)0x1| 0 | |ptr为伪造的全1地址 | 0 | |ptr由__cvta_local_to_generic(0)或__cvta_local_to_generic(1)返回 | 1 | |ptr由__cvta_local_to_generic(全1)返回 | 0 |因此不能仅依据__isLocal(ptr) 1判断ptr是否为可安全访问的栈空间地址。约束说明SIMD与SIMT混合编程场景不支持使用该接口。需要包含的头文件使用该接口需要包含simt_api/device_functions.h头文件。#include simt_api/device_functions.h调用示例SIMT编程场景__global__ __launch_bounds__(1024) void kernel__isLocal(uint32_t* dst, uint32_t* src) { int idx threadIdx.x blockIdx.x * blockDim.x; uint32_t ptr[10]; if(__isLocal(ptr) 1) { dst[idx] 1; } else { dst[idx] 0; } }创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考