clock

产 品 支 持 情 况

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

功 能 说 明

本 接 口 在SIMD、SIMD与SIMT混 合 和SIMT调 试 场 景 中 提 供Clock时 间 戳 功 能，用 于 记 录 从 程 序 启 动 到 接 口 调 用 时 刻 所 经 历 的 时 钟 周 期 数（Cycle Count），便 于 精 确 分 析 执 行 延 迟 和 性 能 瓶 颈。

函 数 原 型

Text

__aicore__ inline uint64_t clock(void)

以 下 接 口 为SIMT VF中 所 使 用 的clock接 口，仅 支 持Ascend 950PR/Ascend 950DT。

Text

__simt_callee__ inline uint64_t clock(void)

参 数 说 明

无

返 回 值 说 明

从 程 序 开 始 到 调 用 时 所 经 历 的 时 钟 周 期 数。

约 束 说 明

无

需 要 包 含 的 头 文 件

使 用 该 接 口 需 要 包 含"utils/debug/asc_time.h"头 文 件。

Text

#include "utils/debug/asc_time.h"

调 用 示 例

• SIMT编 程 场 景：

  Text

  __global__ __launch_bounds__(1024) void SimtKernel(uint64_t* dst)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = clock();
  }
  

• SIMD与SIMT混 合 编 程 场 景：

  Text

  __simt_vf__ __launch_bounds__(1024) inline void SimtKernel(__gm__ uint64_t* dst)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = clock();
  }
  

• SIMD编 程 场 景：

  Text

  __global__ __aicore__ void AicoreKernel(__gm__ uint64_t* dst)
  {
      int idx = AscendC::GetBlockIdx();
      dst[idx] = clock();
  }