pyasc的Python算子生态——用Python语法糖包裹Ascend C的底层能力，为昇腾NPU开发者打开自定义算子的Python大门

前言

在昇腾NPU的算子开发体系中，昇腾CANN一直以Ascend C作为核心编程接口。Ascend C本质上是一种C++扩展语言，开发者用它编写的算子经过毕昇编译器编译后在NPU上执行。这种方案在性能上无可挑剔，但对大量来自深度学习框架生态的开发者而言，存在一个难以回避的现实问题：他们日常使用的语言是Python，而非C++。从PyTorch到NumPy，从数据处理到模型训练，Python已经成为AI开发的事实标准语言。当这些开发者需要在昇腾NPU上写自定义算子时，却被迫切换到C++的思维模式，这种语言鸿沟直接拉高了算子开发的门槛。pyasc项目正是为了弥合这道鸿沟而诞生的。它用原生Python语法封装了Ascend C的全部底层能力，让开发者能够用Python编写在昇腾NPU上运行的自定义算子。这个项目于2025年11月首次上线，目前正处于接口逐步开放的过程中。

为什么Python写算子是开发者更自然的选择

要理解pyasc的价值，需要先看清算子开发者面对的实际处境。自定义算子开发的典型场景包括：研究人员在验证新算法时需要一个高效的前向算子，工程师在部署推理服务时需要把多个操作融合成单个kernel以减少内存搬运开销，或者是框架开发者需要为昇腾NPU适配一些尚未被标准库覆盖的特殊操作。这些场景的共同特征是：开发者需要快速迭代、频繁修改计算逻辑，并且希望在整个流程中尽量减少语言切换带来的认知负担。

用C++写Ascend C算子的痛点是明确的。C++的模板系统、指针操作和内存管理模型，与Python开发者在日常工作中形成的心智模型差异巨大。写一个简单的Add算子，在Ascend C中需要处理GlobalTensor和LocalTensor的模板参数，管理VECIN和VECOUT的逻辑位置，手动编排set_flag和wait_flag的同步时序，还要用<<<…>>>语法来启动核函数。对熟悉C++的系统级开发者来说，这些步骤或许不算复杂，但对以Python为主要工作语言的AI开发者而言，每一层抽象都意味着额外的学习成本和出错概率。

更深层的问题在于开发效率的差距。Python具备动态类型、解释执行、丰富的标准库和极低的代码编辑-运行反馈周期，这些特性让开发者能够把精力集中在计算逻辑本身，而非语言机制上。Python的for循环写起来直觉自然，列表推导式简洁高效，NumPy的广播机制让批量运算变成一行代码。这些语言层面的便利性，在算子开发中体现为更快的原型验证速度和更低的试错成本。

当一个研究者想把论文中的新激活函数在昇腾NPU上跑起来时，他不应该先花几天时间去理解C++模板元编程的细节。pyasc的出发点就是让这些开发者能够用他们最熟悉的语言，直接表达算子的计算意图，同时保证最终产物在NPU上的执行性能与手工编写的Ascend C代码完全一致。

pyasc的设计理念：Pythonic的方式暴露底层硬件能力

pyasc的核心设计原则可以用一句话概括：接口与Ascend C一一对应，语法遵守Python原生规范。这并非简单的"Python绑定"或"封装一层wrapper"，而是一个经过精心设计的编译框架，能够将Python源码通过多级中间表示最终转换为在昇腾NPU上可执行的kernel。

这个转换过程涉及五个核心模块：编译和运行模块、Python前端模块、AST转ASC-IR模块、ASC-IR定义模块、Ascend C代码生成模块。这些模块可以进一步归为前端模块和后端模块两大类。前端包括编译和运行模块、Python前端模块、AST转ASC-IR模块；后端包括ASC-IR定义模块和Ascend C代码生成模块。当开发者在Python中定义一个用@asc.jit装饰的函数时，整个编译链路被自动拉起。

Python前端模块负责提供与Ascend C API一一对应的Python编程接口。这些接口按照Ascend C的分类方式组织为高阶API、基础API、核心数据结构和枚举、框架类接口四大类，分别对应python/asc/language目录下的adv、basic、core、fwk四个子目录。AST转ASC-IR模块遍历Python的抽象语法树，将每个语法节点转换为对应的ASC-IR操作。遍历器的核心代码位于python/asc/codegen/function_visitor.py，它根据不同的AST节点类型调用不同的处理接口。由于ASC-IR的定义和创建接口都在C++侧，这部分通过pybind11机制提供Python调用接口，并使用TableGen工具自动生成绑定定义。

ASC-IR是基于MLIR定义的Dialect，它实现了与Ascend C API的1:1映射。Dialect由Type、Attribute、Interfaces、Operation四个组件构成。Type定义了LocalTensor、GlobalTensor、TQue、TBuf等类型的MLIR表示；Attribute封装了Layout、Format、Event等枚举属性；Interfaces为所有Operation提供统一的API名称获取和注释生成等能力；Operation则是具体的计算操作节点。末尾，Ascend C代码生成模块将ASC-IR翻译为等价的Ascend C代码，再交给毕昇编译器生成NPU上的可执行二进制。

这条链路的关键在于：开发者只看到Python代码，但编译器确保最终生成的机器码与直接用Ascend C编写的版本在语义上完全等价。这意味着pyasc不会带来任何性能损失，Python只是表达层，底层仍然是Ascend C的原生能力在驱动硬件执行计算。

核心语法糖的实现细节

装饰器模式：@asc.jit驱动整个编译流程

pyasc最显眼的语法糖就是@asc.jit装饰器。它承担了多重职责：标记函数需要被JIT编译、接收编译参数、触发整个编译链路。在Ascend C中，核函数的定义需要使用__global__ __aicore__修饰符，而pyasc将其替换为Python开发者更熟悉的装饰器语法。

importascimportasc.runtime.configasconfigimportasc.lib.runtimeasrt TILE_NUM=8BUFFER_NUM=2@asc.jit(kernel_type=config.KernelType.AIC_ONLY)defrelu_pool_kernel(x:asc.GlobalAddress,y:asc.GlobalAddress,pool_size:int,TILE_LEN:asc.ConstExpr[int],BUF_N:asc.ConstExpr[int]):offset=asc.get_block_idx()*pool_size*pool_size x_gm=asc.GlobalTensor()x_gm.set_global_buffer(x+offset,pool_size*pool_size)y_gm=asc.GlobalTensor()y_gm.set_global_buffer(y+offset,pool_size*pool_size)data_type=x.dtype tile_len=pool_size*pool_size//TILE_LEN//BUF_N buf_size=tile_len*BUF_N*data_type.sizeof()x_local=asc.LocalTensor(data_type,asc.TPosition.VECIN,0,tile_len*BUF_N)y_local=asc.LocalTensor(data_type,asc.TPosition.VECOUT,buf_size,tile_len*BUF_N)foriinrange(TILE_LEN*BUF_N):buf=i%BUF_N asc.data_copy(x_local[buf*tile_len:],x_gm[i*tile_len:],tile_len)asc.set_flag(asc.HardEvent.MTE2_V,buf)asc.wait_flag(asc.HardEvent.MTE2_V,buf)asc.relu(y_local[buf*tile_len:],x_local[buf*tile_len:],tile_len)asc.set_flag(asc.HardEvent.V_MTE3,buf)asc.wait_flag(asc.HardEvent.V_MTE3,buf)asc.data_copy(y_gm[i*tile_len:],y_local[buf*tile_len:],tile_len)asc.set_flag(asc.HardEvent.MTE3_MTE2,buf)asc.wait_flag(asc.HardEvent.MTE3_MTE2,buf)

@asc.jit装饰器把Ascend C中__global__ __aicore__的声明式标记转换成了Python生态中最常见的装饰器模式。Python开发者对装饰器非常熟悉——从@property到functools.wraps，装饰器是Python代码组织的基本手段。使用装饰器而非特殊关键字，意味着开发者可以用标准的Python工具链来处理他们的算子代码。装饰器的小括号内传入的kernel_type等编译参数，对应Ascend C中需要通过编译选项或模板参数指定的配置，只是呈现方式更加内聚。ConstExpr泛型标记则解决了Python动态类型与Ascend C编译期常量之间的鸿沟——在Python中变量没有固定类型，ConstExpr让开发者能够显式声明某些参数需要在编译期求值，这个设计在不破坏Python语法的前提下保留了Ascend C对常量表达式的语义要求。编译模块还支持opt_level（毕昇编译器优化级别）、auto_sync（自动同步开关）等参数，这些参数都可以在@asc.jit的括号中传入，保持了配置接口的统一性。

自动类型推导与数据结构映射

在Ascend C中，GlobalTensor和LocalTensor都是模板类，需要显式指定数据类型。Python本身没有静态类型系统，pyasc通过类型注解和运行时推导两种机制来处理这个问题。

importascimporttorchimportasc.lib.runtimeasrt CORE_NUM=8deflaunch_relu_pool(x,pool_size):y=torch.zeros_like(x)block_len=y.numel()//CORE_NUM relu_pool_kernel[CORE_NUM,rt.current_stream()](x,y,block_len)returny

Python的动态类型特性让开发者无需在每个变量声明处重复写数据类型信息。在Ascend C中写AscendC::GlobalTensor这样的模板实例化代码，对C++开发者来说是常规操作，但对Python开发者来说则是额外的认知负担。pyasc在运行时从传入的torch.Tensor中自动提取dtype信息，并将其传递给编译链路。这种方式让算子的定义和调用都保持Python的自然风格。类型注解中的asc.GlobalAddress告诉编译器这是一个全局内存地址参数，而asc.float32这样的类型标记则在需要显式指明类型时提供精确控制。这种"默认自动推导、需要时可显式标注"的策略，既降低了入门门槛，又保留了足够的表达能力。对于LocalTensor的构造，pyasc要求传入dtype、逻辑位置（TPosition.VECIN或TPosition.VECOUT）、起始地址偏移和长度，这四个参数恰好对应Ascend C中分配UB存储时需要指定的全部信息。

核函数调用语法的Python化

Ascend C使用<<<…>>>语法来指定核函数的执行配置（核数、Stream等），这是从CUDA继承来的C++扩展语法。pyasc将其替换为Python中更自然的下标调用方式。

importascimportasc.lib.runtimeasrtimporttorch CORE_NUM=8deflaunch_relu_pool(x:torch.Tensor)->torch.Tensor:y=torch.zeros_like(x)block_len=y.numel()//CORE_NUM relu_pool_kernel[CORE_NUM,rt.current_stream()](x,y,block_len)returny

方括号在Python中有明确的语义——下标访问和切片。pyasc借用这个语法来表达"以某种配置启动kernel"的含义，与Python的__getitem__协议一脉相承。Ascend C中的<<<BLOCK_NUM, STREAM>>>语法在C++中看起来像一个怪异的运算符重载，而pyasc的CORE_NUM, stream则更像是调用一个可配置的函数对象。这种设计让核函数调用看起来和普通Python函数调用几乎一样，只是在函数名后面多了方括号内的运行时配置。开发者不需要学习新的语法，只需要理解方括号内传的是核数和Stream这个约定即可。运行模块会解析方括号内的参数，调用Ascend C Runtime接口加载编译模块生成的kernel二进制并执行。当输入输出张量位于Host侧时，运行模块还会自动完成Host和Device之间的数据拷贝。

Python前端与C++后端的协作模式

理解pyasc的编译架构，有助于开发者判断它的能力边界和适用范围。整个编译流程可以清晰地划分为前端和后端两个半场，两者通过ASC-IR这个统一的中间表示连接。

前端的职责是从Python源码出发，经过解析、转换，生成结构化的中间表示。Python前端模块定义了与Ascend C API一一对应的Python接口，按照功能分模块组织。在基础API目录下，接口又被细分为data_copy.py（数据搬运）、vec_binary.py（双目矢量运算）、vec_unary.py（单目矢量运算）等文件，每个文件对应一类计算操作。这种组织方式与Ascend C的头文件结构保持对应关系。

当JIT编译器被@asc.jit触发时，AST转ASC-IR模块开始遍历Python源码的抽象语法树。pyasc支持的Python语法范围经过仔细筛选：for循环、while循环、条件判断、赋值语句、增强赋值、算术和逻辑运算、下标访问、切片、列表和元组等都在支持之列。这些语法已经足以表达绝大多数算子的计算逻辑。同时，一些在kernel代码中不常见或难以映射到NPU硬件执行的语法被排除在外，例如嵌套函数、lambda、类方法装饰、break/continue、异常处理和标准库调用。

ASC-IR中Operation的命名遵循统一规则：Dialect名加下划线加Ascend C类名加成员函数名。参数映射也有一套明确的顺序：运行时必选参数排在前面，然后是模板必选参数、运行时可选参数、模板可选参数。这套规则确保了Python接口与Ascend C API之间的映射关系清晰可预测。

后端的代码生成模块按API类别分为Basic、Adv、Core、External、Fwk五个子模块，每个子模块负责处理对应类别的ASC-IR到Ascend C的转换规则。生成的Ascend C代码与手工编写的代码在结构上保持一致，因此最终的性能表现也一致。开发者可以通过设置环境变量PYASC_DUMP_PATH来指定输出目录，在编译完成后查看生成的ASC-IR和Ascend C源码，这为调试和性能分析提供了完整的透明度。

JIT编译缓存机制是提升开发体验的关键设计。编译模块会缓存已编译的kernel二进制，避免每次运行都重新编译。缓存的影响因素包括编译选项、kernel参数、全局变量以及被@asc.jit修饰的函数代码本身。开发者可以通过环境变量PYASC_HOME和PYASC_CACHE_DIR控制缓存路径，也可以设置always_compile=True来强制重新编译。在算子开发迭代频繁的场景下，缓存机制能显著缩短代码修改后的验证周期。当开发者修改了kernel函数的实现代码后，缓存会自动失效并重新编译，整个过程对开发者完全透明。

运行模块还支持通过msprof op工具采集profiling数据，生成内存热力图和仿真流水图。这些分析工具帮助开发者定位算子的性能瓶颈，优化数据搬运和计算的重叠比例，调整tile切分策略和双缓冲参数。

实战：用pyasc写一个融合ReLU加池化的预处理算子

图像预处理是计算机视觉推理管线中的常见环节。典型的预处理流程包括ReLU激活和池化操作，这两个操作在昇腾NPU上通常作为独立的算子执行，中间结果需要在Global Memory和Local Memory之间往返搬运。将它们融合为单个kernel可以消除中间数据的搬运开销。下面展示如何用pyasc实现这个融合算子。

完整的算子实现包含核函数和启动函数两部分。核函数用@asc.jit装饰，在单个核上处理分配到的数据切片。启动函数负责准备torch.Tensor、计算数据分块参数，并通过方括号语法指定核数和Stream来调用核函数。

importascimportasc.lib.runtimeasrtimporttorch CORE_NUM=8POOL_SIZE=2TILE_NUM=16BUFFER_NUM=2@asc.jitdeffused_relu_pool_kernel(x:asc.GlobalAddress,y:asc.GlobalAddress,block_len:int,TILE_N:asc.ConstExpr[int],BUF_N:asc.ConstExpr[int],P_SIZE:asc.ConstExpr[int]):offset=asc.get_block_idx()*block_len x_gm=asc.GlobalTensor()x_gm.set_global_buffer(x+offset,block_len)y_gm=asc.GlobalTensor()y_gm.set_global_buffer(y+offset,block_len//(P_SIZE*P_SIZE))dt=x.dtype tile_len=block_len//TILE_N//BUF_N buf_bytes=tile_len*BUF_N*dt.sizeof()x_local=asc.LocalTensor(dt,asc.TPosition.VECIN,0,tile_len*BUF_N)relu_local=asc.LocalTensor(dt,asc.TPosition.VECOUT,buf_bytes,tile_len*BUF_N)y_local=asc.LocalTensor(dt,asc.TPosition.VECOUT,buf_bytes+buf_bytes,tile_len*BUF_N)foriinrange(TILE_N*BUF_N):buf=i%BUF_N src_off=buf*tile_len asc.data_copy(x_local[src_off:],x_gm[i*tile_len:],tile_len)asc.set_flag(asc.HardEvent.MTE2_V,buf)asc.wait_flag(asc.HardEvent.MTE2_V,buf)asc.relu(relu_local[src_off:],x_local[src_off:],tile_len)asc.set_flag(asc.HardEvent.V_MTE3,buf)asc.wait_flag(asc.HardEvent.V_MTE3,buf)asc.data_copy(y_gm[i*tile_len//(P_SIZE*P_SIZE):],y_local[src_off:],tile_len//(P_SIZE*P_SIZE))asc.set_flag(asc.HardEvent.MTE3_MTE2,buf)asc.wait_flag(asc.HardEvent.MTE3_MTE2,buf)deflaunch_fused_relu_pool(x:torch.Tensor)->torch.Tensor:h,w=x.shape[-2],x.shape[-1]out_h,out_w=h//POOL_SIZE,w//POOL_SIZE y=torch.zeros(*x.shape[:-2],out_h,out_w,dtype=x.dtype,device=x.device)block_len=h*w fused_relu_pool_kernel[CORE_NUM,rt.current_stream()](x,y,block_len,TILE_NUM,BUFFER_NUM,POOL_SIZE)returny

这段代码展示了pyasc的几个典型设计特征。ConstExpr标记用于TILE_N、BUF_N、P_SIZE这些需要在编译期确定的参数——它们实际上是常量，在Ascend C中必须作为编译期值传入，而Python没有"编译期常量"的概念，ConstExpr填补了这个语义空白。LocalTensor的构造函数接受dtype、逻辑位置（TPosition.VECIN/VECOUT）、起始地址偏移和长度四个参数，对应Ascend C中需要指定的UB分区信息。整个双缓冲流水线的编排——搬入、ReLU计算、搬出——通过for循环和set_flag/wait_flag完成，与Ascend C中的写法在逻辑结构上完全对应，只是语法换成了Python。这种对应关系让已经了解Ascend C的开发者可以快速上手pyasc，也让pyasc开发者未来可以平滑迁移到Ascend C。三个LocalTensor分别分配在VECIN和VECOUT位置上，地址偏移量通过buf_bytes计算确保不重叠，这是Ascend C编程中管理UB空间的常规手法。

这个融合算子将ReLU和池化合并到同一个kernel中执行，中间结果relu_local始终留在Local Memory中，不需要写回Global Memory再读出来。在单独执行ReLU和池化两个算子的场景下，ReLU的输出必须先搬出到Global Memory，池化算子再从Global Memory搬入。这种额外的内存搬运在昇腾NPU上会消耗额外的MTE带宽和延迟。融合后，数据在Local Memory中直接从ReLU的输出传递到池化的输入，消除了这段搬运路径。启动函数中直接接收torch.Tensor作为输入输出，与PyTorch生态无缝衔接，开发者可以在已有的训练和推理流程中直接调用这个融合kernel。

pyasc与直接编写Ascend C的开发效率对比

用pyasc和直接用Ascend C实现同一个算子，在最终性能上没有差异——编译器保证两者生成等价的kernel二进制。差异体现在开发过程中的各个环节。下面从几个维度进行对比分析。

从上表可以看出，pyasc的核心优势集中在开发效率层面，而非运行时性能层面。它的价值在于降低入门门槛、加快迭代速度、减少语言切换带来的认知开销，同时通过编译链路保证最终产物在性能上与Ascend C手工编写的版本完全等价。对于Python背景的开发者来说，能够用熟悉的语言完成从原型验证到上板运行的全流程，是一种切实的生产力提升。

适用场景分析

pyasc并非要取代Ascend C，而是为算子开发提供了一个更低门槛的入口。它的适用场景有明确的边界。

pyasc最适合的场景是：开发者已经熟悉Python和PyTorch生态，需要在昇腾NPU上实现自定义算子，但不希望投入大量时间学习C++和Ascend C。这在研究团队中尤其常见。研究者可能需要快速验证一个论文中的新算子，或者为特定任务实现一个融合kernel来提升推理吞吐。pyasc支持的Python语法覆盖了算子开发中最常用的控制流和数据操作——for循环、条件判断、赋值、算术运算、下标访问、切片、增强赋值等，这些语法足以表达绝大多数算子的计算逻辑。

pyasc也适合作为Ascend C的学习桥梁。开发者可以先通过pyasc熟悉昇腾NPU的编程模型——Global Memory与Local Memory的数据搬运、双缓冲流水线、事件同步机制——然后再过渡到直接使用Ascend C。pyasc的PYASC_DUMP_PATH机制让开发者可以查看自己Python代码对应的Ascend C输出，这个特性让pyasc成为了理解Ascend C的实践工具。

使用前和使用后的效率对比

总结：

pyasc目前有一些明确的限制需要了解。不支持嵌套函数、lambda表达式、类方法装饰、break和continue语句、异常处理、以及标准库和第三方库函数的调用。这些限制意味着pyasc面向的是相对纯粹的数值计算kernel，而不是通用编程场景。在算子开发中，绝大多数计算逻辑可以用for循环、条件判断和函数调用组合完成，上述不支持的语言特性在kernel代码中本身也很少出现。另外，pyasc支持的接口范围仍在逐步扩展中，当前已开放的接口与Ascend C的完整API集之间还存在差距。如果开发者需要使用尚未被pyasc覆盖的Ascend C接口，仍需直接编写Ascend C代码。

对于需要极致控制硬件行为的场景，或者需要使用pyasc尚未覆盖的Ascend C API的场景，直接编写Ascend C仍然是必要的选择。pyasc的接口与Ascend C一一对应，这意味着它不会提供Ascend C不具备的能力。它的价值是把已有的Ascend C能力用更友好的方式暴露出来，而不是扩展Ascend C本身的功能边界。

仓库地址：https://atomgit.com/cann/pyasc