asc_atomic_or

产 品 支 持 情 况

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

功 能 说 明

对Unified Buffer或Global Memory上address的 数 值 与 指 定 数 值val进 行 原 子 或（|）操 作，即 将address数 值 或（|）val的 结 果 赋 值 到Unified Buffer或Global Memory上。

函 数 原 型

Text

inline int32_t asc_atomic_or(int32_t *address, int32_t val)

Text

inline uint32_t asc_atomic_or(uint32_t *address, uint32_t val)

Text

inline int64_t asc_atomic_or(int64_t *address, int64_t val)

Text

inline uint64_t asc_atomic_or(uint64_t *address, uint64_t val)

参 数 说 明

表 1 参 数 说 明

参 数 名	输 入/输 出	描 述
address	输 出	Unified Buffer或Global Memory的 地 址。
val	输 入	源 操 作 数。

不 同 数 据 类 型 支 持 的 内 存 范 围 说 明 如 下：

表 2 不 同 数 据 类 型 支 持 的 内 存 范 围

参 数 数 据 类 型	支 持 的 内 存 空 间
int32_t、uint32_t	Unified Buffer、Global Memory
int64_t、uint64_t	Global Memory

返 回 值 说 明

Unified Buffer或Global Memory上 的 初 始 数 据。

约 束 说 明

无

需 要 包 含 的 头 文 件

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

Text

#include "simt_api/device_atomic_functions.h"

调 用 示 例

示 例 场 景 为：多 个 线 程 分 别 检 测 到 不 同 特 征，使 用asc_atomic_or接 口 将 这 些 特 征 合 并 到 同 一 个 共 享bitmask中。输 入 参 数 说 明 如 下：

名 称	说 明
observed_flags	每 个 元 素 表 示 一 个 线 程 观 察 到 的 特 征 位。
flags	Global Memory中 的 汇 总bitmask，kernel启 动 前 清 零。
n	特 征 来 源 数 量。

核 心 代 码 实 现 如 下：

• SIMT编 程 场 景：

  C++

  __global__ __launch_bounds__(256) void merge_observed_flags(uint32_t *flags,
                                                             uint32_t *observed_flags,
                                                             uint32_t n)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= n) {
          return;
      }
  
      asc_atomic_or(flags, observed_flags[idx]);
  }
  

• SIMD与SIMT混 合 编 程 场 景：

  SIMD与SIMT混 合 编 程 场 景，需 要 显 式 使 用 地 址 空 间 限 定 符 表 示 地 址 空 间：__gm__表 示Global Memory内 存 空 间，__ubuf__表 示Unified Buffer内 存 空 间。

  C++

  __simt_vf__ __launch_bounds__(1024) inline void merge_observed_flags(__gm__ uint32_t *flags,
                                                                      __gm__ uint32_t *observed_flags,
                                                                      uint32_t n)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= n) {
          return;
      }
  
      asc_atomic_or(flags, observed_flags[idx]);
  }
  

输 出 结 果 示 例 如 下：

Text

observed_flags: 0x1, 0x4, 0x2
flags: 0x7 // 表 明3个 线 程 观 察 到 的 特 征 位 被 合 并