Skip to content

__isGlobal

产品支持情况

  • Ascend 950PR/Ascend 950DT:支持
  • Atlas A3 训练系列产品/Atlas A3 推理系列产品:不支持
  • Atlas A2 训练系列产品/Atlas A2 推理系列产品:不支持
  • Atlas 200I/500 A2 推理产品:不支持
  • Atlas 推理系列产品AI Core:不支持
  • Atlas 推理系列产品Vector Core:不支持
  • Atlas 训练系列产品:不支持

功能说明

判断输入的指针是否指向Global Memory内存空间的地址。

函数原型

C++
unsigned int __isGlobal(const void* ptr)

参数说明

表1 参数说明

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

返回值说明

如果输入的指针指向Global Memory内存空间的地址,则返回1,否则返回0。
该接口根据输入指针的地址空间信息进行分类判断,不校验该指针是否为可安全访问的有效地址。__isGlobal返回1仅表示该指针被分类为Global Memory地址,不代表该地址一定可以安全访问。特殊场景说明如下:

输入场景返回值
ptr为有效Unified Buffer指针0
ptr为有效栈空间指针0
ptrnullptr0
ptr(void*)0x10
ptr为全1地址1
ptr__cvta_global_to_generic(0)__cvta_global_to_generic(1)返回0
ptr__cvta_global_to_generic(全1)返回1

因此,不能仅依据__isGlobal(ptr) == 1判断ptr是否为可安全访问的Global Memory地址。

约束说明

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

需要包含的头文件

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

C++
#include "simt_api/device_functions.h"

调用示例

SIMT编程场景:

C++
__global__ __launch_bounds__(1024) void kernel_isGlobal(uint32_t* dst, uint32_t* src)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = __isGlobal(src + idx);
}

免责声明:本站内容由 asc-devkit 仓 master 分支自动编译生成,属于持续开发版本,可能存在缺陷,仅供预览与参考。如需稳定及商用资料,请查阅官方 昇腾社区