cat 迁移说明

1. 算子说明

• 算子名称：torch.cat
• 迁移模式：torch_npu
• 来源与交付形态摘要：来源为 PyTorch Shape.cu 内建 CUDA backend；交付为 ported-ops/cat 下的 out-of-tree torch_npu PrivateUse1 扩展。
• 功能概述：将多个同形状兼容 Tensor 沿指定维度拼接，支持 torch.cat(tensors, dim) 与 torch.cat(tensors, dim, out=out)。
• 输入输出说明：输入为 dense strided NPU Tensor 列表；输出为同设备 NPU Tensor。
• 一比一复刻结论：保留 CUDA 版两级 parallel path、四个 copy kernel、dtype 分发、32-bit index gate、contiguous / non-contiguous gate 与 narrow + copy_ fallback；未接受单 kernel 或 host fallback 降级。

2. 原始 CUDA 实现摘要

• 原始实现文件：/home/y00621698/simt-buddy/pytorch/aten/src/ATen/native/cuda/Shape.cu
• Kernel 入口：CatArrayBatchedCopy、CatArrayBatchedCopy_contig、CatArrayBatchedCopy_vectorized、CatArrayBatchedCopy_alignedK_contig
• 用户侧调用路径摘要：torch.cat / torch.cat(..., out=...) 经 aten::cat / aten::cat.out dispatch 到 CUDA cat_out_cuda。

3. Ascend C SIMT 迁移摘要

3.1 迁移策略

• 使用 torch_npu PrivateUse1 扩展注册 aten::cat 与 aten::cat.out。
• 将 CUDA __device__ helper 改为 __aicore__，kernel launch 保留 CUDA-like SIMT 语法。
• 将 CUDA stream、device property、vector memory helper 替换为 NPU / ACL / SIMT counterpart。
• 是否实现可复用 Ascend counterpart：是，提供 SimtVec、AlignedVector、NPU stream launch wrapper 和 cat metadata dispatch。

3.2 主要改动与结论

• 新增 simt_cat/csrc/simt/cat_simt.asc：包含 precompute、host dispatch 和四个 SIMT kernel。
• 新增 simt_cat/csrc/cat_bindings.asc：注册 aten::cat / aten::cat.out。
• 新增 setup.py、pyproject.toml、requirements.txt：可安装扩展工程。
• 新增 test/test_cat.py：Python 调用验证。
• 重大降级决策（如有）：无。用户预设不同意降级，本次未合并 kernel、未缩 dtype、未删除 fallback。

3.3 API 与语法替换说明

CUDA 项	Ascend 项	说明
__device__	__aicore__	device helper qualifier 替换
at::cuda::getCurrentCUDAStream()	c10_npu::getCurrentNPUStream().stream(true)	NPU stream
multiProcessorCount	ACL_DEV_ATTR_VECTOR_CORE_NUM	grid 计算硬件并行度来源
C10_CUDA_KERNEL_LAUNCH_CHECK	OpCommand::RunOpApiV2	torch_npu SIMT launch 包装
MemoryAccess.cuh vector helper	SimtVec / AlignedVector	本地 SIMT vector counterpart

3.4 依赖与能力覆盖结论

• 依赖闭包处理：ATen 公共 shape / dtype / resize helper 复用；CUDAContext 和 CUDA memory helper 迁移为 NPU/SIMT counterpart。
• 能力覆盖结论：dense strided cat 主路径已迁移；5 个 Python 行为用例通过。
• shape / kernel 选择行为保留情况：保留 CAT_ARRAY_BATCH_SIZE=128、non-contiguous 64/64 batch、4D 限制、32-bit indexing gate、vectorized/alignedK/contig/general/fallback 分支。
• 语法规避依据：__device__ 按 grammar.md 改为 __aicore__；dim3 显式初始化适配 Ascend SIMT 编译器。
• 显式降级项：无用户接受的降级项。

4. 目录结构

ported-ops/cat/
├── plan.md
├── README.md
├── pyproject.toml
├── requirements.txt
├── setup.py
├── simt_cat/
│   ├── __init__.py
│   ├── _C.cpython-313-x86_64-linux-gnu.so
│   └── csrc/
│       ├── cat_bindings.asc
│       ├── cat_simt.h
│       └── simt/
│           └── cat_simt.asc
└── test/
    └── test_cat.py

5. 构建方式

5.1 环境准备

source /usr/local/Ascend/cann/set_env.sh

5.2 编译步骤

cd /home/y00621698/simt-buddy/tasks/ported-ops/cat
python setup.py build_ext --inplace

默认使用 -O0，因为当前 bisheng -O3 在 aligned/vector 模板路径触发后端错误。可用以下命令复现阻塞：

ASCEND_OPT_FLAG=-O3 python setup.py build_ext --inplace

5.3 安装步骤

cd /home/y00621698/simt-buddy/tasks/ported-ops/cat
python -m pip install -e . --no-build-isolation

5.4 出包步骤

cd /home/y00621698/simt-buddy/tasks/ported-ops/cat
python -m pip wheel . --no-build-isolation

6. 验证与使用方式

6.1 原生/C++ 侧验证

cd /home/y00621698/simt-buddy/tasks/ported-ops/cat
python setup.py build_ext --inplace

6.2 Python 侧验证

cd /home/y00621698/simt-buddy/tasks/ported-ops/cat
python test/test_cat.py -v

6.3 出包验证

cd /home/y00621698/simt-buddy/tasks/ported-ops/cat
python -m pip install -e . --no-build-isolation
python test/test_cat.py -v

使用时导入扩展后直接调用 torch.cat：

import torch
import torch_npu
import simt_cat

x = torch.arange(12, dtype=torch.float32).reshape(3, 4).npu()
y = torch.arange(8, dtype=torch.float32).reshape(2, 4).npu()
z = torch.cat([x, y], dim=0)

7. 最终验证结果摘要

• 验证环境：torch 2.11.0+cpu、torch_npu 2.11.0.dev20260414、CANN /usr/local/Ascend/cann-9.0.0、bisheng 可用。
• Ascend 950 PR 实机验证：未完整确认；当前环境 torch.npu.is_available() == True 且测试在 NPU device 0 运行，但缺少 npu-smi 等工具确认板卡型号。
• 构建结果：python setup.py build_ext --inplace 通过。
• 原生/C++ 侧验证结果：生成 _C.cpython-313-x86_64-linux-gnu.so。
• Python 侧验证结果：python test/test_cat.py -v，6 个用例中 5 个通过，1 个 channels-last 因 torch_npu allocator 限制跳过。
• 出包/安装结果：python -m pip install -e . --no-build-isolation 成功，simt_cat-0.0.1 installed。
• 结果校验结论：已执行 contiguous、float16、non-contiguous、out dtype cast、single-input fallback 的 CPU reference 对比，结果一致。
• 当前状态标签：migrated but not validated on Ascend 950 PR

8. 已知限制

• 显式降级项：无。
• 已阻塞项：channels-last 验证被当前 torch_npu allocator 阻塞，错误为 Only c10::MemoryFormat::Contiguous is supported for creating a npu tensor。
• 用户已接受的降级项：无。
• 其他限制：默认 -O0 构建；-O3 当前触发 bisheng 后端错误。

9. 问题排查记录

• 问题阶段：release 构建。
• 问题现象：ASCEND_OPT_FLAG=-O3 python setup.py build_ext --inplace 报 Copy one register into another with a different width。
• 根因判断：bisheng 后端优化阶段处理 aligned/vector 模板路径失败。
• 已验证动作：改用 -O0 后构建通过；未删除 vectorized/alignedK 路径。
• 相关语法/约束依据：本地 constraints 未列出该后端优化问题，按 build blocker 记录。

10. 后续建议

• 使用可确认 Ascend 950 PR 型号的环境重新运行 python test/test_cat.py -v。
• torch_npu 支持 channels-last 创建后，去掉 skip 并验证 channels-last 分支。
• bisheng 修复 -O3 后端问题后，以 ASCEND_OPT_FLAG=-O3 重跑构建与测试。

11. 未完成事项

• Ascend 950 PR 型号确认。
• channels-last 分支结果校验。
• -O3 构建验证。