__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 |
ptr为nullptr | 0 |
ptr为(void*)0x1 | 0 |
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);
}