Star732
1031
代码
代码
Issues54
Pull Requests176
Actions
项目讨论
Wiki
分析
Star732
1031
已合并
fix example #5449
watermelon_milk创建于 10 天前
fix example #5449
已合并
watermelon_milk创建于 10 天前
从watermelon_milk/asc-devkit_newtensor_snake合入到cann/asc-devkitmaster
共 7 个文件变更+33-35
M.pre-commit-config.yaml+1-1
@@ -58,7 +58,7 @@ repos:
58 - id: codespell58 - id: codespell
59 args: [59 args: [
60 "-L",60 "-L",
61- "CANN,cann,NNAL,nnal,ASCEND,ascend,EnQue,CopyIn,ArchType,AND,ND,tbe,copyin,alog,Collet,Enque,Normale,OUTTER,OffsetT,REGIST,Te,ThirdParty,VOR,abl,ans,catched,commiter,devlop,dout,enque,ges,nIn,nd,outter,padd,parm,subtile",61+ "CANN,cann,NNAL,nnal,ASCEND,ascend,EnQue,CopyIn,ArchType,AND,ND,tbe,copyin,alog,Collet,Enque,Normale,OUTTER,OffsetT,REGIST,Te,ThirdParty,VOR,abl,ans,catched,commiter,devlop,dout,enque,ges,nIn,nd,outter,padd,parm,subtile,te",
62 "--skip",62 "--skip",
63 "*.py,*.cpp,*.hpp,*.c,*.h,*.asc,*.cc",63 "*.py,*.cpp,*.hpp,*.c,*.h,*.asc,*.cc",
64 ]64 ]
Mdocs/en/asc_how_to_choose_api.md+1-1
@@ -121,7 +121,7 @@ C++-level low-level programming interface for in-depth development of high-perfo
121- Production-grade high-performance operator development scenarios that require guaranteed code standardization, maintainability and scalability while pursuing ultimate hardware performance.121- Production-grade high-performance operator development scenarios that require guaranteed code standardization, maintainability and scalability while pursuing ultimate hardware performance.
122 122 
123**Reference Examples**123**Reference Examples**
124-- [Matmul Example Based on Tensor API](../../examples/01_simd_cpp_api/03_basic_api/03_matrix_compute/mmad_tensor_api/mmad_tensor_api.asc)124+- [Matmul Example Based on Tensor API](../../examples/01_simd_cpp_api/07_tensor_api/mmad_tensor_api/mmad_tensor_api.asc)
125 125 
126---126---
127 127 
Mdocs/zh/asc_how_to_choose_api.md+1-1
@@ -131,7 +131,7 @@ graph TD
131 131 
132**参考示例**132**参考示例**
133 133 
134-- [基于Tensor API的Matmul示例](../../examples/01_simd_cpp_api/03_basic_api/03_matrix_compute/mmad_tensor_api/mmad_tensor_api.asc)134+- [基于Tensor API的Matmul示例](../../examples/01_simd_cpp_api/07_tensor_api/mmad_tensor_api/mmad_tensor_api.asc)
135 135 
136### 基础API（Basic API）：轻量化Tensor编程接口（TPipe/TQue自动化资源管理）136### 基础API（Basic API）：轻量化Tensor编程接口（TPipe/TQue自动化资源管理）
137 137 
Mexamples/01_simd_cpp_api/07_tensor_api/matmul_mxfp4_tensor_api_high_performance/README.md+1-1
@@ -239,7 +239,7 @@ for nBlock in N blocks:
239计算阶段的关键参数包括 `m/n/k` 尺寸、是否初始化 C 矩阵：239计算阶段的关键参数包括 `m/n/k` 尺寸、是否初始化 C 矩阵：
240 240 
241```cpp241```cpp
242-asc::te::mmad_params mmadParams{curM, curN, curK, unit_flag_mode::disable, true};242+asc::te::mmad_params mmadParams{curM, curN, curK, asc::te::unit_flag_mode::disable, true};
243mmadParams.init_with_zero = (kBlockIdx == 0);243mmadParams.init_with_zero = (kBlockIdx == 0);
244asc::te::mmad(mmadAtom.with(mmadParams), l0TensorC, l0TensorA, l0TensorB);244asc::te::mmad(mmadAtom.with(mmadParams), l0TensorC, l0TensorA, l0TensorB);
245```245```
Mexamples/01_simd_cpp_api/07_tensor_api/matmul_mxfp4_tensor_api_high_performance/README_en.md+1-1
@@ -239,7 +239,7 @@ for nBlock in N blocks:
239Key parameters in the computation stage include `m/n/k` dimensions and whether to initialize the C matrix:239Key parameters in the computation stage include `m/n/k` dimensions and whether to initialize the C matrix:
240 240 
241```cpp241```cpp
242-asc::te::mmad_params mmadParams{curM, curN, curK, unit_flag_mode::disable, true};242+asc::te::mmad_params mmadParams{curM, curN, curK, asc::te::unit_flag_mode::disable, true};
243mmadParams.init_with_zero = (kBlockIdx == 0);243mmadParams.init_with_zero = (kBlockIdx == 0);
244asc::te::mmad(mmadAtom.with(mmadParams), l0TensorC, l0TensorA, l0TensorB);244asc::te::mmad(mmadAtom.with(mmadParams), l0TensorC, l0TensorA, l0TensorB);
245```245```
Mexamples/01_simd_cpp_api/07_tensor_api/matmul_mxfp4_tensor_api_high_performance/mmad_mx.asc+14-19
@@ -49,8 +49,7 @@ struct KernelTrait {
49 static constexpr uint32_t scaleFactorK = scaleFactorK_;49 static constexpr uint32_t scaleFactorK = scaleFactorK_;
50};50};
51 51 
52-constexpr asc::te::mmad_trait MX_MMAD_TRAIT =52+constexpr asc::te::mmad_trait MX_MMAD_TRAIT = asc::te::mmad_trait{0, false, false, true, asc::te::mmad_type::mx};
53- asc::te::mmad_trait{0, false, false, true, asc::te::mmad_type::mx};
54struct MmadTraitMX {53struct MmadTraitMX {
55 using TraitType = asc::te::mmad_trait;54 using TraitType = asc::te::mmad_trait;
56 static constexpr const TraitType value = MX_MMAD_TRAIT;55 static constexpr const TraitType value = MX_MMAD_TRAIT;
@@ -74,11 +73,9 @@ public:
74 asc::te::make_mem_ptr(b), asc::te::make_frame_layout<asc::te::dn_ext_layout_ptn>(Trait::K, Trait::N));73 asc::te::make_mem_ptr(b), asc::te::make_frame_layout<asc::te::dn_ext_layout_ptn>(Trait::K, Trait::N));
75 74 
76 auto gmTensorAs = asc::te::make_tensor(75 auto gmTensorAs = asc::te::make_tensor(
77- asc::te::make_mem_ptr(as),76+ asc::te::make_mem_ptr(as), asc::te::make_frame_layout<asc::te::scalea_nd_layout_ptn>(Trait::M, scaleK));
78- asc::te::make_frame_layout<asc::te::scalea_nd_layout_ptn>(Trait::M, scaleK));
79 auto gmTensorBs = asc::te::make_tensor(77 auto gmTensorBs = asc::te::make_tensor(
80- asc::te::make_mem_ptr(bs),78+ asc::te::make_mem_ptr(bs), asc::te::make_frame_layout<asc::te::scaleb_dn_layout_ptn>(scaleK, Trait::N));
81- asc::te::make_frame_layout<asc::te::scaleb_dn_layout_ptn>(scaleK, Trait::N));
82 auto gmTensorC = asc::te::make_tensor(79 auto gmTensorC = asc::te::make_tensor(
83 asc::te::make_mem_ptr(c), asc::te::make_frame_layout<asc::te::nd_ext_layout_ptn>(Trait::M, Trait::N));80 asc::te::make_mem_ptr(c), asc::te::make_frame_layout<asc::te::nd_ext_layout_ptn>(Trait::M, Trait::N));
84 81 
@@ -90,11 +87,9 @@ public:
90 asc::te::make_coord(0, nIterIdx * Trait::singleCoreN),87 asc::te::make_coord(0, nIterIdx * Trait::singleCoreN),
91 asc::te::make_shape(Trait::singleCoreK, actualSingleCoreN));88 asc::te::make_shape(Trait::singleCoreK, actualSingleCoreN));
92 auto gmSingleTensorAs = gmTensorAs.slice(89 auto gmSingleTensorAs = gmTensorAs.slice(
93- asc::te::make_coord(mIterIdx * Trait::singleCoreM, 0),90+ asc::te::make_coord(mIterIdx * Trait::singleCoreM, 0), asc::te::make_shape(actualSingleCoreM, scaleK));
94- asc::te::make_shape(actualSingleCoreM, scaleK));
95 auto gmSingleTensorBs = gmTensorBs.slice(91 auto gmSingleTensorBs = gmTensorBs.slice(
96- asc::te::make_coord(0, nIterIdx * Trait::singleCoreN),92+ asc::te::make_coord(0, nIterIdx * Trait::singleCoreN), asc::te::make_shape(scaleK, actualSingleCoreN));
97- asc::te::make_shape(scaleK, actualSingleCoreN));
98 auto gmSingleTensorC = gmTensorC.slice(93 auto gmSingleTensorC = gmTensorC.slice(
99 asc::te::make_coord(mIterIdx * Trait::singleCoreM, nIterIdx * Trait::singleCoreN),94 asc::te::make_coord(mIterIdx * Trait::singleCoreM, nIterIdx * Trait::singleCoreN),
100 asc::te::make_shape(actualSingleCoreM, actualSingleCoreN));95 asc::te::make_shape(actualSingleCoreM, actualSingleCoreN));
@@ -147,14 +142,14 @@ private:
147 auto l0PtrAPong = asc::te::make_mem_ptr(l0BufAPong);142 auto l0PtrAPong = asc::te::make_mem_ptr(l0BufAPong);
148 auto l0PtrBPing = asc::te::make_mem_ptr(l0BufBPing);143 auto l0PtrBPing = asc::te::make_mem_ptr(l0BufBPing);
149 auto l0PtrBPong = asc::te::make_mem_ptr(l0BufBPong);144 auto l0PtrBPong = asc::te::make_mem_ptr(l0BufBPong);
150- auto l0PtrAsPing = asc::te::make_mem_ptr<asc::te::location::l0scalea, fp8_e8m0_t>(145+ auto l0PtrAsPing =
151- reinterpret_cast<uint64_t>(l0BufAPing) / 16);146+ asc::te::make_mem_ptr<asc::te::location::l0scalea, fp8_e8m0_t>(reinterpret_cast<uint64_t>(l0BufAPing) / 16);
152- auto l0PtrAsPong = asc::te::make_mem_ptr<asc::te::location::l0scalea, fp8_e8m0_t>(147+ auto l0PtrAsPong =
153- reinterpret_cast<uint64_t>(l0BufAPong) / 16);148+ asc::te::make_mem_ptr<asc::te::location::l0scalea, fp8_e8m0_t>(reinterpret_cast<uint64_t>(l0BufAPong) / 16);
154- auto l0PtrBsPing = asc::te::make_mem_ptr<asc::te::location::l0scaleb, fp8_e8m0_t>(149+ auto l0PtrBsPing =
155- reinterpret_cast<uint64_t>(l0BufBPing) / 16);150+ asc::te::make_mem_ptr<asc::te::location::l0scaleb, fp8_e8m0_t>(reinterpret_cast<uint64_t>(l0BufBPing) / 16);
156- auto l0PtrBsPong = asc::te::make_mem_ptr<asc::te::location::l0scaleb, fp8_e8m0_t>(151+ auto l0PtrBsPong =
157- reinterpret_cast<uint64_t>(l0BufBPong) / 16);152+ asc::te::make_mem_ptr<asc::te::location::l0scaleb, fp8_e8m0_t>(reinterpret_cast<uint64_t>(l0BufBPong) / 16);
158 153 
159 auto l1TensorAPing = asc::te::make_tensor(asc::te::make_mem_ptr(l1BufAPing), l1LayoutA);154 auto l1TensorAPing = asc::te::make_tensor(asc::te::make_mem_ptr(l1BufAPing), l1LayoutA);
160 auto l1TensorAPong = asc::te::make_tensor(asc::te::make_mem_ptr(l1BufAPong), l1LayoutA);155 auto l1TensorAPong = asc::te::make_tensor(asc::te::make_mem_ptr(l1BufAPong), l1LayoutA);
@@ -323,7 +318,7 @@ private:
323 318 
324 // ---- Compute: Mmad matrix multiply-accumulate ----319 // ---- Compute: Mmad matrix multiply-accumulate ----
325 asc_lock(PIPE_M, l0Slot);320 asc_lock(PIPE_M, l0Slot);
326- asc::te::mmad_params mmadParams{curM, curN, curK, unit_flag_mode::disable, true};321+ asc::te::mmad_params mmadParams{curM, curN, curK, asc::te::unit_flag_mode::disable, true};
327 mmadParams.init_with_zero = (kBlockIdx == 0);322 mmadParams.init_with_zero = (kBlockIdx == 0);
328 323 
329 asc::te::mmad(mmadAtom.with(mmadParams), l0TensorC, l0TensorA, l0TensorB);324 asc::te::mmad(mmadAtom.with(mmadParams), l0TensorC, l0TensorA, l0TensorB);
Mexamples/01_simd_cpp_api/07_tensor_api/mmad_tensor_api/mmad_tensor_api.asc+14-11
@@ -61,14 +61,16 @@ __cube__ __global__ void mmad_quant_custom(
61 auto l0BPtr = make_mem_ptr(l0BBuf);61 auto l0BPtr = make_mem_ptr(l0BBuf);
62 auto l0CPtr = make_mem_ptr(l0CBuf);62 auto l0CPtr = make_mem_ptr(l0CBuf);
63 63 
64- auto copyGM2L1Atom = make_copy(copy_gm_to_l1{}, copy_gm_to_l1_trait_default{});64+ auto copyGM2L1Atom = make_copy(copy_gm_to_l1{}, gm_to_l1_trait_default{});
65- auto copyL12L0AAtom = make_copy(copy_l1_to_l0a{}, copy_l1_to_l0a_trait_default{});65+ auto copyL12L0AAtom = make_copy(copy_l1_to_l0a{}, l1_to_l0a_trait_default{});
66- auto copyL12L0BAtom = make_copy(copy_l1_to_l0b{}, copy_l1_to_l0b_trait_default{});66+ auto copyL12L0BAtom = make_copy(copy_l1_to_l0b{}, l1_to_l0b_trait_default{});
67- auto copyL0C2GMAtom = make_copy(copy_l0c_to_gm{}, copy_l0c_to_gm_trait_default{});67+ auto copyL0C2GMAtom = make_copy(copy_l0c_to_gm{}, l0c_to_gm_trait_default{});
68 auto mmadAtom = make_mmad(mmad_operation{}, mmad_trait_default{});68 auto mmadAtom = make_mmad(mmad_operation{}, mmad_trait_default{});
69 69 
70- auto l1ATensor = make_tensor(make_mem_ptr(l1ABuf), make_frame_layout<nz_layout_ptn, A_TYPE>(baseM * stepM, baseK * stepK));70+ auto l1ATensor =
71- auto l1BTensor = make_tensor(make_mem_ptr(l1BBuf), make_frame_layout<zn_layout_ptn, B_TYPE>(baseK * stepK, baseN * stepN));71+ make_tensor(make_mem_ptr(l1ABuf), make_frame_layout<nz_layout_ptn, A_TYPE>(baseM * stepM, baseK * stepK));
72+ auto l1BTensor =
73+ make_tensor(make_mem_ptr(l1BBuf), make_frame_layout<zn_layout_ptn, B_TYPE>(baseK * stepK, baseN * stepN));
72 74 
73 uint32_t mLoop = (actualSingleCoreM + baseM - 1) / baseM;75 uint32_t mLoop = (actualSingleCoreM + baseM - 1) / baseM;
74 uint32_t nLoop = (actualSingleCoreN + baseN - 1) / baseN;76 uint32_t nLoop = (actualSingleCoreN + baseN - 1) / baseN;
@@ -130,9 +132,9 @@ __cube__ __global__ void mmad_quant_custom(
130 para.m = curM;132 para.m = curM;
131 para.n = curN;133 para.n = curN;
132 para.k = curK;134 para.k = curK;
133- para.cmatrix_init_val = (kIter == 0) ? true : false;135+ para.init_with_zero = (kIter == 0) ? true : false;
134 if constexpr (ENABLE_UNITFLAG) {136 if constexpr (ENABLE_UNITFLAG) {
135- para.unit_flag = (kIter == kLoop - 1) ? 3 : 2;137+ para.unit_flag = (kIter == kLoop - 1) ? unit_flag_mode::enable_update : unit_flag_mode::enable_keep;
136 }138 }
137 mmad(mmadAtom.with(para), l0CTensor, l0ATensor, l0BTensor);139 mmad(mmadAtom.with(para), l0CTensor, l0ATensor, l0BTensor);
138 if (kIter == kLoop - 1) {140 if (kIter == kLoop - 1) {
@@ -145,14 +147,15 @@ __cube__ __global__ void mmad_quant_custom(
145 asc_lock(PIPE_MTE2, QUANT_Index);147 asc_lock(PIPE_MTE2, QUANT_Index);
146 auto gmCSlice = gmCSingleTensor.slice(make_coord(i * baseM, j * baseN), make_shape(curM, curN));148 auto gmCSlice = gmCSingleTensor.slice(make_coord(i * baseM, j * baseN), make_shape(curM, curN));
147 auto l1QuantTensor = make_tensor(l1QuantPtr, make_frame_layout<nd_ext_layout_ptn>(1, curN));149 auto l1QuantTensor = make_tensor(l1QuantPtr, make_frame_layout<nd_ext_layout_ptn>(1, curN));
148- copy(copyGM2L1Atom, l1QuantTensor, gmQuantSingleTensor.slice(make_coord(0, j * baseN), make_shape(1, curN)));150+ copy(
151+ copyGM2L1Atom, l1QuantTensor, gmQuantSingleTensor.slice(make_coord(0, j * baseN), make_shape(1, curN)));
149 asc_unlock(PIPE_MTE2, QUANT_Index);152 asc_unlock(PIPE_MTE2, QUANT_Index);
150 asc_lock(PIPE_FIX, QUANT_Index);153 asc_lock(PIPE_FIX, QUANT_Index);
151 154 
152 asc_lock(PIPE_FIX, L0C_Index);155 asc_lock(PIPE_FIX, L0C_Index);
153 if constexpr (ENABLE_UNITFLAG) {156 if constexpr (ENABLE_UNITFLAG) {
154- fixpipe_params params;157+ l0c_to_gm_params params;
155- params.unit_flag = 3;158+ params.unit_flag = unit_flag_mode::enable_update;
156 copy(copyL0C2GMAtom.with(params), gmCSlice, l0CTensor, l1QuantTensor);159 copy(copyL0C2GMAtom.with(params), gmCSlice, l0CTensor, l1QuantTensor);
157 } else {160 } else {
158 copy(copyL0C2GMAtom, gmCSlice, l0CTensor, l1QuantTensor);161 copy(copyL0C2GMAtom, gmCSlice, l0CTensor, l1QuantTensor);