asc_atomic_dec

产 品 支 持 情 况

• 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的 数 值 进 行 原 子 减1操 作，如 果address上 的 数 值 等 于0或 大 于 指 定 数 值val，则 对address赋 值 为val，否 则 将address上 数 值 减1。

函 数 原 型

Text

inline uint32_t asc_atomic_dec(uint32_t *address, uint32_t val)

Text

inline uint64_t asc_atomic_dec(uint64_t *address, uint64_t val)

参 数 说 明

表 1 参 数 说 明

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

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

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

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

返 回 值 说 明

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

约 束 说 明

无

需 要 包 含 的 头 文 件

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

Text

#include "simt_api/device_atomic_functions.h"

调 用 示 例

示 例 场 景 为：多 个 线 程 从 高 到 低 循 环 分 配 槽 位，使 用asc_atomic_dec接 口 获 取 更 新 前 的 旧 计 数。当 旧 值 为0时，计 数 器 会 回 绕 到 指 定 上 界capacity - 1。输 入 参 数 说 明 如 下：

名 称	说 明
ticket	Global Memory中 的 反 向 环 形 计 数 器，kernel启 动 前 初 始 化。
slots	保 存 每 个 线 程 获 得 的 槽 位 编 号。
capacity	环 形 队 列 容 量。
n	需 要 分 配 槽 位 的 线 程 数。

核 心 代 码 实 现 如 下：

• SIMT编 程 场 景：

  C++

  __global__ __launch_bounds__(256) void allocate_reverse_ring_slot(uint32_t *ticket,
                                                                  uint32_t *slots,
                                                                  uint32_t capacity,
                                                                  uint32_t n)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= n) {
          return;
      }
  
      uint32_t old_ticket = asc_atomic_dec(ticket, capacity - 1U);
      slots[idx] = old_ticket;
  }
  

• SIMD与SIMT混 合 编 程 场 景：

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

  C++

  __simt_vf__ __launch_bounds__(1024) inline void allocate_reverse_ring_slot(__gm__ uint32_t *ticket,
                                                                           __gm__ uint32_t *slots,
                                                                           uint32_t capacity,
                                                                           uint32_t n)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= n) {
          return;
      }
  
      uint32_t old_ticket = asc_atomic_dec(ticket, capacity - 1U);
      slots[idx] = old_ticket;
  }
  

输 出 结 果 示 例 如 下：

C++

ticket before: 0
capacity: 4
n: 6
slots: 0, 3, 2, 1, 0, 3 // 顺 序 由 实 际 原 子 执 行 顺 序 决 定
ticket after: 2