| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
simt算子风格改为c风格 Co-authored-by: chenfeng61<chenfeng61@huawei.com> # message auto-generated for no-merge-commit merge: !4243 merge simt_zg into master simt算子风格改为c风格 Created-by: chenfeng61 Commit-by: chenfeng61 Merged-by: cann-robot Description: ## 描述 将 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(...) | ### 新增头文件依赖 cpp #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修复 - [ ] 新特性 - [x] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述：SIMT API 迁移/重构 See merge request: cann/ops-nn!4243 | 3 个月前 | |
nn仓新增deformable_offsets_grad算子实现 Co-authored-by: xxxzzzh<xiongzhihao4@huawei.com> # message auto-generated for no-merge-commit merge: !2357 merge add_deform_offsets_grad into master nn仓新增deformable_offsets_grad算子实现 Created-by: xxxzzzh Commit-by: xxxzzzh Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> nn仓新增deformable_offsets_grad算子实现，原来存在于ops-cv仓，用来补齐DeformableConv2dBackward算子能力(deformable_offsets_grad+Conv2dBackward) ## 算子功能： DeformableOffsets算子的目的是根据offsets（kernel采样点的偏移值）来收集用于卷积的特征采样点，并对其进行重组，方便Conv2d算子进行卷积计算。而DeformableOffsetsGrad即为这一过程的反向。 ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #000--> https://gitcode.com/cann/ops-nn/issues/1332 ## 测试 <!--描述进行了哪些测试来验证你的改动。--> deformable_offsets_grad算子UT通过，examples执行通过 ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> 新增了deformable_offsets_grad算子README ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/ops-nn!2357 | 5 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 3 个月前 | ||
| 5 个月前 |