asc_atomic_min

产 品 支 持 情 况

• 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数 据 做 原 子 求 最 小 值 操 作，即 将Unified Buffer或Global Memory的 数 据 与 指 定 数 据 中 的 最 小 值 赋 值 到Unified Buffer或Global Memory地 址 中。

函 数 原 型

Text

inline int32_t asc_atomic_min(int32_t *address, int32_t val)

Text

inline uint32_t asc_atomic_min(uint32_t *address, uint32_t val)

Text

inline float asc_atomic_min(float *address, float val)

Text

inline int64_t asc_atomic_min(int64_t *address, int64_t val)

Text

inline uint64_t asc_atomic_min(uint64_t *address, uint64_t val)

Text

inline half asc_atomic_min(half *address, half val)

Text

inline bfloat16_t asc_atomic_min(bfloat16_t *address, bfloat16_t val)

Text

inline half2 asc_atomic_min(half2 *address, half2 val)

Text

inline bfloat16x2_t asc_atomic_min(bfloat16x2_t *address, bfloat16x2_t val)

参 数 说 明

表 1 参 数 说 明

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

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

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

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

返 回 值 说 明

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

注 意，由 于 底 层 硬 件 约 束，half和bfloat16_t类 型 的 返 回 值 不 准 确，避 免 直 接 使 用 这 些 类 型 的 返 回 值。

约 束 说 明

无

需 要 包 含 的 头 文 件

使 用 除half、half2、bfloat16_t、bfloat16x2_t类 型 之 外 的 接 口 需 要 包 含"simt_api/device_atomic_functions.h"头 文 件，使 用half和half2类 型 接 口 需 要 包 含"simt_api/asc_fp16.h"头 文 件，使 用bfloat16_t和bfloat16x2_t类 型 接 口 需 要 包 含"simt_api/asc_bf16.h"头 文 件。

Text

#include "simt_api/device_atomic_functions.h"

Text

#include "simt_api/asc_fp16.h"

Text

#include "simt_api/asc_bf16.h"

调 用 示 例

示 例 场 景 为：多 个 线 程 扫 描 延 迟 数 组，使 用asc_atomic_min接 口 将 全 局 最 低 延 迟 写 入 同 一 个 结 果 地 址。输 入 参 数 说 明 如 下：

名 称	说 明
latency	每 个 元 素 表 示 一 次 请 求 的 延 迟 值。
min_latency	Global Memory中 的 最 小 值 结 果，kernel启 动 前 初 始 化 为 足 够 大 的 值。
n	延 迟 样 本 数 量。

核 心 代 码 实 现 如 下：

• SIMT编 程 场 景：

  C++

  __global__ __launch_bounds__(256) void find_min_latency(uint32_t *min_latency,
                                                         uint32_t *latency,
                                                         uint32_t n)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= n) {
          return;
      }
  
      asc_atomic_min(min_latency, latency[idx]);
  }
  

• SIMD与SIMT混 合 编 程 场 景：

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

  C++

  __simt_vf__ __launch_bounds__(1024) inline void find_min_latency(__gm__ uint32_t *min_latency,
                                                                  __gm__ uint32_t *latency,
                                                                  uint32_t n)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= n) {
          return;
      }
  
      asc_atomic_min(min_latency, latency[idx]);
  }
  

输 出 结 果 示 例 如 下：

Text

latency: 31, 20, 45
min_latency: 20 // 表 明 所 有 线 程 并 发 更 新 后 得 到 最 小 值