__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;
    }
}