simt算子风格改为c风格-ops-nn-AtomGit

已合并

simt算子风格改为c风格 #4243

chenfeng创建于 4月25日从chenfeng61/ops-nn： simt_zg 合入到 cann/ops-nn： master

simt算子风格改为c风格 #4243

已合并

chenfeng创建于 4月25日

从chenfeng61/ops-nn simt_zg合入到cann/ops-nn master

chenfeng成员

4月25日

描述

将 arch35 下 conv/loss 算子 kernel 中的 SIMT API 从旧的c++风格改为c风格

改动范围

涉及 8 个算子, 11 个 kernel 文件:

算子	文件
deformable_offsets	`conv/deformable_offsets/op_kernel/arch35/deformable_offsets.h`
deformable_offsets_grad	`conv/deformable_offsets_grad/op_kernel/arch35/deformable_offsets_grad.h`
ctc_loss_v2	`loss/ctc_loss_v2/op_kernel/arch35/ctc_loss_v2_base.h`
ctc_loss_v2	`loss/ctc_loss_v2/op_kernel/arch35/ctc_loss_v2_fp16.h`
ctc_loss_v2	`loss/ctc_loss_v2/op_kernel/arch35/ctc_loss_v2_fp32.h`
ctc_loss_v2_grad	`loss/ctc_loss_v2_grad/op_kernel/arch35/ctc_loss_v2_grad.h`
nll_loss	`loss/nll_loss/op_kernel/arch35/nll_loss_mix.h`
nll_loss	`loss/nll_loss/op_kernel/arch35/nll_loss_simt.h`
nll_loss_grad	`loss/nll_loss_grad/op_kernel/arch35/nll_loss_grad.h`
sparse_softmax_cross_entropy_with_logits	`loss/sparse_softmax_cross_entropy_with_logits/op_kernel/arch35/..._full_load.h`
sparse_softmax_cross_entropy_with_logits	`loss/sparse_softmax_cross_entropy_with_logits/op_kernel/arch35/..._split_r.h`

API 映射关系

旧 API (`Simt::`)	新 API
`Simt::GetThreadIdx()` / `Simt::GetThreadIdx<0>()`	`threadIdx.x`
`Simt::GetThreadIdx<1>()`	`threadIdx.y`
`Simt::GetThreadNum()` / `Simt::GetThreadNum<0>()`	`blockDim.x`
`Simt::GetThreadNum<1>()`	`blockDim.y`
`Simt::GetBlockIdx()`	`blockIdx.x`
`Simt::GetBlockNum()`	`gridDim.x`
`Simt::ThreadBarrier()`	`asc_syncthreads()`
`Simt::VF_CALL<...>(Simt::Dim3{...}, ...)`	`asc_vf_call<...>(dim3{...}, ...)`
`Simt::AtomicAdd(...)`	`asc_atomic_add(...)`
`Simt::Log1p(...)`	`log1pf(...)`
`Simt::Exp(...)`	`expf(...)`
`Simt::Max(...)`	`fmaxf(...)`

新增头文件依赖

#include "simt_api/asc_simt.h"          // asc_vf_call / asc_syncthreads / dim3
#include "simt_api/device_atomic_functions.h"  // asc_atomic_add
#include "simt_api/asc_fp16.h"          // FP16 SIMT 支持
#include "simt_api/asc_bf16.h"          // BF16 SIMT 支持

关联的Issue

https://gitcode.com/cann/ops-nn/issues/2670

测试

本地编译通过验证
算子精度测试

文档更新

无

类型标签

Bug修复
新特性
性能优化
文档更新
其他，请描述：SIMT API 迁移/重构

Pull Request已成功合入, 合并人@CANN-robot

（感谢 chenfeng 的贡献）

chenfeng成员

4月25日创建了 pull request，commit 8aef4ad1

CANN-robot成员

4月25日评论：

Thanks for your pull-request.
The full list of commands accepted by me can be found at here。
You can get sig-info at here

PR Approval Progress

✅ Congratulations! All modules have met the lgtm and approve requirements.

Module Approval Details

module	lgtm status	approve status
conv	✅ 余鑫, 张瑜翔, 范其瑞, 周奇龙 (4/2)	✅ 余鑫, 张瑜翔 (2/1)
loss	✅ 周奇龙, 范其瑞 (2/2)	✅ 周奇龙 (1/1)

💡 Tip:

Committer can comment /approve or /lgtm

Commenting /approve implies both code review (lgtm) and intent to merge (approve)

CLA Signature Pass

chenfeng61, thanks for your pull request. All authors of the commits have signed the CLA. 👍

CANN-robot成员

4月25日添加了label：cann-cla/yes

CANN-robot成员

4月25日将crystalhu,yu-xinjie62,yangyang016,fanqirui,zhou-qilong,chenqi317,lileizheng,liubo75,tangweiwei2,zhangyuxiang0119,wangcimeng,wangyongguang,zhang-wu设为评审人

CANN-robot成员

4月25日将crystalhu,yu-xinjie62,yangyang016,fanqirui,zhou-qilong,chenqi317,lileizheng,liubo75,tangweiwei2,zhangyuxiang0119,wangcimeng,wangyongguang,zhang-wu设为审查人

此处折叠了64条消息查看更多

sunday成员

5月13日评论：

/lgtm
/approve

CANN-robot成员

5月13日添加了label：lgtmapproved

CANN-robot成员

5月13日评论：

Review Guide

This pull-request passes review.
Committers who wrote a comment of /approve are: 余鑫, 张瑜翔, 周奇龙.
Reviewers who wrote a comment of /lgtm are: 余鑫, 张瑜翔, 范其瑞, 周奇龙.

CANN-robot成员

5月13日关闭了关联的issue

CANN-robot成员

5月13日合入了pull request