已合并
保持third_party下目录结构和ops-tensor一致 #324
Crrryyyy创建于 6月30日
保持third_party下目录结构和ops-tensor一致 #324
已合并
共 45 个文件变更+111-177
| @@ -184,7 +184,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 184 | para.k = curKL0; | 184 | para.k = curKL0; |
| 185 | 185 | ||
| 186 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 186 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 187 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 187 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 188 | 188 | ||
| 189 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(tool::ZERO_FLAG); | 189 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(tool::ZERO_FLAG); |
| 190 | } | 190 | } |
| @@ -319,7 +319,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 319 | 319 | ||
| 320 | // Perform MAD (Multiply-Add) operation: C += A * B | 320 | // Perform MAD (Multiply-Add) operation: C += A * B |
| 321 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 321 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 322 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 322 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 323 | 323 | ||
| 324 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 324 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 325 | l0PingPong_++; // Toggle L0 buffer | 325 | l0PingPong_++; // Toggle L0 buffer |
| @@ -212,7 +212,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 212 | para.k = curKL0; | 212 | para.k = curKL0; |
| 213 | 213 | ||
| 214 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 214 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 215 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 215 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 216 | 216 | ||
| 217 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 217 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 218 | l0PingPong++; // Toggle L0 buffer | 218 | l0PingPong++; // Toggle L0 buffer |
| @@ -217,7 +217,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 217 | } | 217 | } |
| 218 | 218 | ||
| 219 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 219 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 220 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 220 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 221 | 221 | ||
| 222 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 222 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 223 | l0PingPong_++; // Toggle L0 buffer | 223 | l0PingPong_++; // Toggle L0 buffer |
| @@ -227,7 +227,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 227 | } | 227 | } |
| 228 | 228 | ||
| 229 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 229 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 230 | - AscendC::Te::Copy(copyL0C2GM, tensorCGmBlock, tensorL0C, AscendC::Te::FixpipeParams{tool::FINAL_ACCUMULATION}); | 230 | + AscendC::Te::Copy(copyL0C2GM.with(AscendC::Te::FixpipeParams{tool::FINAL_ACCUMULATION}), tensorCGmBlock, tensorL0C); |
| 231 | } | 231 | } |
| 232 | 232 | ||
| 233 | // Final synchronization waits | 233 | // Final synchronization waits |
| @@ -204,14 +204,14 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 204 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 204 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 205 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 205 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 206 | 206 | ||
| 207 | - AscendC::MmadParams para; | 207 | + AscendC::Te::MmadParams para; |
| 208 | para.cmatrixInitVal = (iter1 == 0 && iter0 == 0); | 208 | para.cmatrixInitVal = (iter1 == 0 && iter0 == 0); |
| 209 | para.m = curM; | 209 | para.m = curM; |
| 210 | para.n = curN; | 210 | para.n = curN; |
| 211 | para.k = curKL0; | 211 | para.k = curKL0; |
| 212 | 212 | ||
| 213 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 213 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 214 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 214 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 215 | 215 | ||
| 216 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 216 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 217 | l0PingPong++; | 217 | l0PingPong++; |
| @@ -213,7 +213,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 213 | para.k = curKL0; | 213 | para.k = curKL0; |
| 214 | 214 | ||
| 215 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 215 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 216 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 216 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 217 | 217 | ||
| 218 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 218 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 219 | l0PingPong++; | 219 | l0PingPong++; |
| @@ -207,7 +207,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 207 | para.k = curKL0; | 207 | para.k = curKL0; |
| 208 | 208 | ||
| 209 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 209 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 210 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 210 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 211 | 211 | ||
| 212 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 212 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 213 | l0PingPong_++; | 213 | l0PingPong_++; |
| @@ -398,8 +398,8 @@ __global__ __aicore__ void MatmulKernel( | |||
| 398 | auto tensorScaleAL0 = | 398 | auto tensorScaleAL0 = |
| 399 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); | 399 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); |
| 400 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510Atom{}); | 400 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510Atom{}); |
| 401 | - AscendC::Te::Copy( | 401 | + CopyL12L0MxScaleA.Call( |
| 402 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | 402 | + tensorScaleAL0, tensorBlockScaleAL1, |
| 403 | AscendC::Te::MakeCoord(0, scaleKL1IterOffset + kL0Offset)); | 403 | AscendC::Te::MakeCoord(0, scaleKL1IterOffset + kL0Offset)); |
| 404 | 404 | ||
| 405 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<tool::C0_SIZE>>(curKL0, curN); | 405 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<tool::C0_SIZE>>(curKL0, curN); |
| @@ -412,8 +412,8 @@ __global__ __aicore__ void MatmulKernel( | |||
| 412 | auto tensorScaleBL0 = | 412 | auto tensorScaleBL0 = |
| 413 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); | 413 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); |
| 414 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510Atom{}); | 414 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510Atom{}); |
| 415 | - AscendC::Te::Copy( | 415 | + CopyL12L0MxScaleB.Call( |
| 416 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | 416 | + tensorScaleBL0, tensorBlockScaleBL1, |
| 417 | AscendC::Te::MakeCoord(scaleKL1IterOffset + kL0Offset, 0)); | 417 | AscendC::Te::MakeCoord(scaleKL1IterOffset + kL0Offset, 0)); |
| 418 | 418 | ||
| 419 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 419 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| @@ -451,7 +451,7 @@ __global__ __aicore__ void MatmulKernel( | |||
| 451 | AscendC::WaitFlag<AscendC::HardEvent::M_FIX>(tool::ZERO_FLAG); | 451 | AscendC::WaitFlag<AscendC::HardEvent::M_FIX>(tool::ZERO_FLAG); |
| 452 | 452 | ||
| 453 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 453 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 454 | - AscendC::Te::Copy(CopyL0C2GM, tensorCGmBlock, tensorL0C, AscendC::Te::FixpipeParams{tool::FINAL_ACCUMULATION}); | 454 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{tool::FINAL_ACCUMULATION}), tensorCGmBlock, tensorL0C); |
| 455 | 455 | ||
| 456 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(tool::ZERO_FLAG); | 456 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(tool::ZERO_FLAG); |
| 457 | } | 457 | } |
| @@ -226,7 +226,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 226 | para.k = curKL0; | 226 | para.k = curKL0; |
| 227 | 227 | ||
| 228 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 228 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 229 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 229 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 230 | 230 | ||
| 231 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 231 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 232 | l0PingPong_++; | 232 | l0PingPong_++; |
| @@ -56,7 +56,7 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 56 | // -------------------- AIV 逻辑（AI Vector Core） -------------------- | 56 | // -------------------- AIV 逻辑（AI Vector Core） -------------------- |
| 57 | if ASCEND_IS_AIV { | 57 | if ASCEND_IS_AIV { |
| 58 | // 判断当前AIV核心是否已完成所有预期的循环轮次（lastLoopTotalCnt * 任务分配比） | 58 | // 判断当前AIV核心是否已完成所有预期的循环轮次（lastLoopTotalCnt * 任务分配比） |
| 59 | - if (curBlockIdx >= lastLoopTotalCnt * AscendC::GetTaskRation()) { | 59 | + if (curBlockIdx >= lastLoopTotalCnt * AscendC::GetTaskRatio()) { |
| 60 | AscendC::CrossCoreWaitFlag<tool::AIC_SYNC_AIV_MODE_4, PIPE_MTE3>(tool::AIC_SYNC_AIV_FLAG); | 60 | AscendC::CrossCoreWaitFlag<tool::AIC_SYNC_AIV_MODE_4, PIPE_MTE3>(tool::AIC_SYNC_AIV_FLAG); |
| 61 | AscendC::SyncAll(); | 61 | AscendC::SyncAll(); |
| 62 | return; | 62 | return; |
| @@ -244,12 +244,12 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 244 | 244 | ||
| 245 | // newBlockIdx: 重新映射后的块索引（对应 M-N 平面上的块编号） | 245 | // newBlockIdx: 重新映射后的块索引（对应 M-N 平面上的块编号） |
| 246 | // 计算方式：当前核心索引 curBlockIdx 除以 (任务分配比 × K轴分片数) | 246 | // 计算方式：当前核心索引 curBlockIdx 除以 (任务分配比 × K轴分片数) |
| 247 | - // AscendC::GetTaskRation() 获取任务分配比率（AIV 核心数 / AIC 核心数） | 247 | + // AscendC::GetTaskRatio() 获取任务分配比率（AIV 核心数 / AIC 核心数） |
| 248 | - uint64_t newBlockIdx = curBlockIdx / (AscendC::GetTaskRation() * skKTileNum); | 248 | + uint64_t newBlockIdx = curBlockIdx / (AscendC::GetTaskRatio() * skKTileNum); |
| 249 | 249 | ||
| 250 | // kTileIdx: 重新映射后的 K 维度分块索引 | 250 | // kTileIdx: 重新映射后的 K 维度分块索引 |
| 251 | // 通过对 (任务分配比 × K轴分片数) 取模得到 | 251 | // 通过对 (任务分配比 × K轴分片数) 取模得到 |
| 252 | - uint64_t kTileIdx = curBlockIdx % (AscendC::GetTaskRation() * skKTileNum); | 252 | + uint64_t kTileIdx = curBlockIdx % (AscendC::GetTaskRatio() * skKTileNum); |
| 253 | uint64_t cGmIndex = newBlockIdx + (mTileNum * nTileNum - (mTileNum * nTileNum) % blockNum); | 253 | uint64_t cGmIndex = newBlockIdx + (mTileNum * nTileNum - (mTileNum * nTileNum) % blockNum); |
| 254 | uint64_t mTileIdx = cGmIndex / nTileNum; | 254 | uint64_t mTileIdx = cGmIndex / nTileNum; |
| 255 | uint64_t nTileIdx = cGmIndex % nTileNum; | 255 | uint64_t nTileIdx = cGmIndex % nTileNum; |
| @@ -367,4 +367,4 @@ python3 profile_matmul.py 1024 2048 1024 | |||
| 367 | 367 | ||
| 368 | ## 6. 支持架构 | 368 | ## 6. 支持架构 |
| 369 | 369 | ||
| 370 | -NPU ARCH 3510 | 370 | +NPU ARCH 3510 |
| @@ -324,7 +324,7 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 324 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 324 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 325 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 325 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 326 | 326 | ||
| 327 | - AscendC::MmadParams para; | 327 | + AscendC::Te::MmadParams para; |
| 328 | para.cmatrixInitVal = (iter1 == 0 && iter0 == 0); | 328 | para.cmatrixInitVal = (iter1 == 0 && iter0 == 0); |
| 329 | para.m = curM; | 329 | para.m = curM; |
| 330 | para.n = curN; | 330 | para.n = curN; |
| @@ -337,7 +337,7 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 337 | } | 337 | } |
| 338 | 338 | ||
| 339 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 339 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 340 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 340 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 341 | 341 | ||
| 342 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 342 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 343 | l0PingPong++; // Toggle L0 buffer | 343 | l0PingPong++; // Toggle L0 buffer |
| @@ -345,10 +345,9 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 345 | if (iter0 + 1 == kL1TileNum) { | 345 | if (iter0 + 1 == kL1TileNum) { |
| 346 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 346 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 347 | if (tool::CheckIsSkScene(tmpTileIdx, blockNum, tileNum) && tailMNTileNum <= (blockNum / tool::NUM_TWO)) { | 347 | if (tool::CheckIsSkScene(tmpTileIdx, blockNum, tileNum) && tailMNTileNum <= (blockNum / tool::NUM_TWO)) { |
| 348 | - AscendC::Te::Copy( | 348 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams(tool::FINAL_ACCUMULATION)), gmWorkSpace, tensorL0C); |
| 349 | - CopyL0C2GM, gmWorkSpace, tensorL0C, AscendC::Te::FixpipeParams(tool::FINAL_ACCUMULATION)); | ||
| 350 | } else { | 349 | } else { |
| 351 | - AscendC::Te::Copy(CopyL0C2GM, tensorCGmBlock, tensorL0C, AscendC::Te::FixpipeParams(tool::FINAL_ACCUMULATION)); | 350 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams(tool::FINAL_ACCUMULATION)), tensorCGmBlock, tensorL0C); |
| 352 | } | 351 | } |
| 353 | } | 352 | } |
| 354 | AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(l1BufId); | 353 | AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(l1BufId); |
| @@ -373,7 +372,7 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 373 | 372 | ||
| 374 | // Idle cores exit early after synchronization | 373 | // Idle cores exit early after synchronization |
| 375 | uint64_t curBlockIdxInAiv = AscendC::GetBlockIdx(); | 374 | uint64_t curBlockIdxInAiv = AscendC::GetBlockIdx(); |
| 376 | - if (curBlockIdxInAiv >= lastLoopTotalCnt * AscendC::GetTaskRation()) { | 375 | + if (curBlockIdxInAiv >= lastLoopTotalCnt * AscendC::GetTaskRatio()) { |
Y | |||
| 377 | AscendC::CrossCoreWaitFlag<tool::AIC_SYNC_AIV_MODE_4, PIPE_MTE3>(tool::AIC_SYNC_AIV_FLAG); | 376 | AscendC::CrossCoreWaitFlag<tool::AIC_SYNC_AIV_MODE_4, PIPE_MTE3>(tool::AIC_SYNC_AIV_FLAG); |
| 378 | AscendC::SyncAll(); | 377 | AscendC::SyncAll(); |
| 379 | return; | 378 | return; |
| @@ -384,9 +383,9 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 384 | AscendC::SyncAll(); | 383 | AscendC::SyncAll(); |
| 385 | 384 | ||
| 386 | // Decompose balanced workload index back to (M,N) tile and K-tile | 385 | // Decompose balanced workload index back to (M,N) tile and K-tile |
| 387 | - uint64_t newBlockIdx = AscendC::GetBlockIdx() / (AscendC::GetTaskRation() * skKTileNum); | 386 | + uint64_t newBlockIdx = AscendC::GetBlockIdx() / (AscendC::GetTaskRatio() * skKTileNum); |
| 388 | uint64_t cGmIndex = newBlockIdx + (mTileNum * nTileNum - (mTileNum * nTileNum) % blockNum); | 387 | uint64_t cGmIndex = newBlockIdx + (mTileNum * nTileNum - (mTileNum * nTileNum) % blockNum); |
| 389 | - uint64_t kTileIdx = AscendC::GetBlockIdx() % (AscendC::GetTaskRation() * skKTileNum); | 388 | + uint64_t kTileIdx = AscendC::GetBlockIdx() % (AscendC::GetTaskRatio() * skKTileNum); |
| 390 | uint64_t mTileIdx = cGmIndex / nTileNum; | 389 | uint64_t mTileIdx = cGmIndex / nTileNum; |
| 391 | uint64_t nTileIdx = cGmIndex % nTileNum; | 390 | uint64_t nTileIdx = cGmIndex % nTileNum; |
| 392 | 391 | ||
| @@ -398,11 +397,11 @@ __global__ __aicore__ __mix__(1, 2) void MatmulKernel( | |||
| 398 | AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(tool::ZERO_FLAG); | 397 | AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(tool::ZERO_FLAG); |
| 399 | } | 398 | } |
| 400 | 399 | ||
| 401 | - uint64_t aivMte2Num = tool::CheckIsSkScene(0, blockNum, tileNum) ? AscendC::GetTaskRation() : AscendC::BLOCK_CUBE; | 400 | + uint64_t aivMte2Num = tool::CheckIsSkScene(0, blockNum, tileNum) ? AscendC::GetTaskRatio() : AscendC::BLOCK_CUBE; |
| 402 | 401 | ||
| 403 | for (uint64_t index = 0; index < aivMte2Num; ++index) { | 402 | for (uint64_t index = 0; index < aivMte2Num; ++index) { |
| 404 | uint64_t mBurstBase = | 403 | uint64_t mBurstBase = |
| 405 | - tool::CeilAlign(tool::CeilDiv(curM, skKTileNum * AscendC::GetTaskRation()), tool::CeilDiv(tool::BLOCK_BYTE_SIZE, curN)); | 404 | + tool::CeilAlign(tool::CeilDiv(curM, skKTileNum * AscendC::GetTaskRatio()), tool::CeilDiv(tool::BLOCK_BYTE_SIZE, curN)); |
| 406 | uint64_t mBurstCnt = tool::CeilDiv(curM, mBurstBase); | 405 | uint64_t mBurstCnt = tool::CeilDiv(curM, mBurstBase); |
| 407 | uint64_t mBurstTail = curM - (mBurstCnt - 1) * mBurstBase; | 406 | uint64_t mBurstTail = curM - (mBurstCnt - 1) * mBurstBase; |
| 408 | if (kTileIdx >= mBurstCnt) { | 407 | if (kTileIdx >= mBurstCnt) { |
| @@ -274,7 +274,7 @@ __global__ __aicore__ void MatmulKernel(GM_ADDR aGm, GM_ADDR bGm, GM_ADDR cGm, u | |||
| 274 | para.k = curKL0; | 274 | para.k = curKL0; |
| 275 | 275 | ||
| 276 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); | 276 | auto MadOp = AscendC::Te::MakeMmad(AscendC::Te::MmadOperation{}, AscendC::Te::MmadTraitDefault{}); |
| 277 | - AscendC::Te::Mmad(MadOp, tensorL0C, tensorAL0, tensorBL0, para); | 277 | + AscendC::Te::Mmad(MadOp.with(para), tensorL0C, tensorAL0, tensorBL0); |
| 278 | 278 | ||
| 279 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); | 279 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0BufId); |
| 280 | l0PingPong++; | 280 | l0PingPong++; |
MSamples/2_Performance/full_quant_fused_infer_attention_score_story/include/block/flash_attention_score_block_vec_infer.h+1-1
| @@ -577,7 +577,7 @@ __aicore__ inline void FABlockVecInfer<TEMPLATE_ARGS>::InitOutputSingleCore(Cons | |||
| 577 | TEMPLATES_DEF_NO_DEFAULT | 577 | TEMPLATES_DEF_NO_DEFAULT |
| 578 | __aicore__ inline void FABlockVecInfer<TEMPLATE_ARGS>::InitLseOutputSingleCore(ConstInfo<isInfer, hasRope> &constInfo) | 578 | __aicore__ inline void FABlockVecInfer<TEMPLATE_ARGS>::InitLseOutputSingleCore(ConstInfo<isInfer, hasRope> &constInfo) |
| 579 | { | 579 | { |
| 580 | - int64_t coreNum = GetBlockNum() * GetTaskRation(); | 580 | + int64_t coreNum = GetBlockNum() * GetTaskRatio(); |
| 581 | auto &initParams = this->tilingData->initOutputParams; | 581 | auto &initParams = this->tilingData->initOutputParams; |
| 582 | if (coreNum != 0 && constInfo.aivIdx < coreNum) { | 582 | if (coreNum != 0 && constInfo.aivIdx < coreNum) { |
| 583 | int64_t singleCoreLseSize = initParams.totalSoftMaxLseOutputSize / coreNum; | 583 | int64_t singleCoreLseSize = initParams.totalSoftMaxLseOutputSize / coreNum; |
MSamples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/include/block/quant_grouped_matmul_hif8_block_mmad.h+4-10
| @@ -177,13 +177,9 @@ public: | |||
| 177 | 177 | ||
| 178 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 178 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 179 | if constexpr (isScalarScale) { | 179 | if constexpr (isScalarScale) { |
| 180 | - AscendC::Te::Copy( | 180 | + AscendC::Te::Copy(copyL0C2GM.with(AscendC::Te::FixpipeParams{GroupedMatmulRecipe::FINAL_ACCUMULATION}), gmC, c1Local, scalarScale_); |
| 181 | - copyL0C2GM, gmC, c1Local, scalarScale_, | ||
| 182 | - AscendC::Te::FixpipeParams{GroupedMatmulRecipe::FINAL_ACCUMULATION}); | ||
| 183 | } else { | 181 | } else { |
| 184 | - AscendC::Te::Copy( | 182 | + AscendC::Te::Copy(copyL0C2GM.with(AscendC::Te::FixpipeParams{GroupedMatmulRecipe::FINAL_ACCUMULATION}), gmC, c1Local, tensorX2L1); |
| 185 | - copyL0C2GM, gmC, c1Local, tensorX2L1, | ||
| 186 | - AscendC::Te::FixpipeParams{GroupedMatmulRecipe::FINAL_ACCUMULATION}); | ||
| 187 | } | 183 | } |
| 188 | 184 | ||
| 189 | if (enableL0cPingPong_) { | 185 | if (enableL0cPingPong_) { |
| @@ -239,10 +235,8 @@ private: | |||
| 239 | params.cmatrixInitVal = | 235 | params.cmatrixInitVal = |
| 240 | (initMode == L0_INIT_MODE_ABL1) ? (initQ0 == 0 && iter1 == 0) | 236 | (initMode == L0_INIT_MODE_ABL1) ? (initQ0 == 0 && iter1 == 0) |
| 241 | : (initQ0 == 0 && initQ1 == 0 && iter1 == 0); | 237 | : (initQ0 == 0 && initQ1 == 0 && iter1 == 0); |
| 242 | - AscendC::Te::Mmad( | 238 | + AscendC::Te::Mmad(AscendC::Te::MmadAtom< |
| 243 | - AscendC::Te::MmadAtom< | 239 | + AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, AscendC::Te::MmadTraitDefault>>{}.with(params), c1Local, l0aLocal, l0bLocal); |
| 244 | - AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, AscendC::Te::MmadTraitDefault>>{}, | ||
| 245 | - c1Local, l0aLocal, l0bLocal, params); | ||
| 246 | 240 | ||
| 247 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0PingPong_ & 0x1); | 241 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0PingPong_ & 0x1); |
| 248 | l0PingPong_++; | 242 | l0PingPong_++; |
MSamples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/include/block/quant_grouped_matmul_mx_block_mmad.h+3-8
| @@ -368,17 +368,13 @@ private: | |||
| 368 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), | 368 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), |
| 369 | AscendC::Te::MakeFrameLayout<AscendC::Te::ZZLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 369 | AscendC::Te::MakeFrameLayout<AscendC::Te::ZZLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| 370 | curM, CeilDiv(curKL0, GroupedMatmulRecipe::MX_DIVISOR_SIZE) * GroupedMatmulRecipe::MX_MULTI_SIZE)); | 370 | curM, CeilDiv(curKL0, GroupedMatmulRecipe::MX_DIVISOR_SIZE) * GroupedMatmulRecipe::MX_MULTI_SIZE)); |
| 371 | - AscendC::Te::Copy( | 371 | + CopyL12L0MxScaleA3510.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kaL1Offset + kL0Offset)); |
| 372 | - CopyL12L0MxScaleA3510, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 373 | - AscendC::Te::MakeCoord(0, kaL1Offset + kL0Offset)); | ||
| 374 | 372 | ||
| 375 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( | 373 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( |
| 376 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), | 374 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), |
| 377 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 375 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| 378 | CeilDiv(curKL0, GroupedMatmulRecipe::MX_DIVISOR_SIZE) * GroupedMatmulRecipe::MX_MULTI_SIZE, curN)); | 376 | CeilDiv(curKL0, GroupedMatmulRecipe::MX_DIVISOR_SIZE) * GroupedMatmulRecipe::MX_MULTI_SIZE, curN)); |
| 379 | - AscendC::Te::Copy( | 377 | + CopyL12L0MxScaleB3510.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kbL1Offset + kL0Offset, 0)); |
| 380 | - CopyL12L0MxScaleB3510, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 381 | - AscendC::Te::MakeCoord(kbL1Offset + kL0Offset, 0)); | ||
| 382 | 378 | ||
| 383 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 379 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 384 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 380 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -488,8 +484,7 @@ private: | |||
| 488 | 484 | ||
| 489 | // Single fixpipe for this block's C tile (bf16 out); L0C ping-pong only when enabled by tiling. | 485 | // Single fixpipe for this block's C tile (bf16 out); L0C ping-pong only when enabled by tiling. |
| 490 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 486 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 491 | - AscendC::Te::Copy( | 487 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{GroupedMatmulRecipe::FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 492 | - CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{GroupedMatmulRecipe::FINAL_ACCUMULATION}); | ||
| 493 | if (enableL0cPingPong_) { | 488 | if (enableL0cPingPong_) { |
| 494 | l0cPingPong_++; | 489 | l0cPingPong_++; |
| 495 | } | 490 | } |
MSamples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/include/block/quant_grouped_matmul_mx_block_scheduler_split_m.h+1-1
| @@ -51,7 +51,7 @@ private: | |||
| 51 | int32_t mBaseTail_; | 51 | int32_t mBaseTail_; |
| 52 | int32_t nBaseTail_; | 52 | int32_t nBaseTail_; |
| 53 | uint32_t blockNum_ = AscendC::GetBlockNum(); | 53 | uint32_t blockNum_ = AscendC::GetBlockNum(); |
| 54 | - uint32_t blockIdx_ = AscendC::GetBlockIdx() / AscendC::GetTaskRation(); | 54 | + uint32_t blockIdx_ = AscendC::GetBlockIdx() / AscendC::GetTaskRatio(); |
| 55 | uint32_t startBlockIdx_; | 55 | uint32_t startBlockIdx_; |
| 56 | uint32_t endBlockIdx_{blockNum_ - 1}; | 56 | uint32_t endBlockIdx_{blockNum_ - 1}; |
| 57 | 57 | ||
MSamples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/include/block/weight_quant_grouped_matmul_mxfp8fp4_block_mmad_split_m.h+3-7
| @@ -250,9 +250,7 @@ __aicore__ inline void WQBMM_CUBE_COMPUTE_CLASS::ProcessTileL1( | |||
| 250 | auto tensorScaleAL0 = | 250 | auto tensorScaleAL0 = |
| 251 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0BufIdx_ * L0_BUF_OFFSET), layoutScaleAL0); | 251 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0BufIdx_ * L0_BUF_OFFSET), layoutScaleAL0); |
| 252 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(Tile::CopyL12L0MxScaleA3510{}); | 252 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(Tile::CopyL12L0MxScaleA3510{}); |
| 253 | - AscendC::Te::Copy( | 253 | + CopyL12L0MxScaleA3510.Call(tensorScaleAL0, tensorScaleAL1_, AscendC::Te::MakeCoord(0, ((l1KOffset + kbOffset) % MX_SCALE_K_L1_SIZE))); |
| 254 | - CopyL12L0MxScaleA3510, tensorScaleAL0, tensorScaleAL1_, | ||
| 255 | - AscendC::Te::MakeCoord(0, ((l1KOffset + kbOffset) % MX_SCALE_K_L1_SIZE))); | ||
| 256 | 254 | ||
| 257 | auto CopyL12L0B = AscendC::Te::MakeCopy(AscendC::Te::CopyL12L0B{}); | 255 | auto CopyL12L0B = AscendC::Te::MakeCopy(AscendC::Te::CopyL12L0B{}); |
| 258 | auto layoutBL0 = MakeLayoutBL0{}(realL0k, param.nL1Size); | 256 | auto layoutBL0 = MakeLayoutBL0{}(realL0k, param.nL1Size); |
| @@ -266,9 +264,7 @@ __aicore__ inline void WQBMM_CUBE_COMPUTE_CLASS::ProcessTileL1( | |||
| 266 | auto tensorScaleBL0 = | 264 | auto tensorScaleBL0 = |
| 267 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0BufIdx_ * L0_BUF_OFFSET), layoutScaleBL0); | 265 | AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0BufIdx_ * L0_BUF_OFFSET), layoutScaleBL0); |
| 268 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(Tile::CopyL12L0MxScaleB3510{}); | 266 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(Tile::CopyL12L0MxScaleB3510{}); |
| 269 | - AscendC::Te::Copy( | 267 | + CopyL12L0MxScaleB3510.Call(tensorScaleBL0, tensorScaleBL1_, AscendC::Te::MakeCoord(((l1KOffset + kbOffset) % MX_SCALE_K_L1_SIZE), 0)); |
| 270 | - CopyL12L0MxScaleB3510, tensorScaleBL0, tensorScaleBL1_, | ||
| 271 | - AscendC::Te::MakeCoord(((l1KOffset + kbOffset) % MX_SCALE_K_L1_SIZE), 0)); | ||
| 272 | 268 | ||
| 273 | SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM_); | 269 | SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM_); |
| 274 | WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM_); | 270 | WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM_); |
| @@ -364,7 +360,7 @@ __aicore__ inline void WQBMM_CUBE_COMPUTE_CLASS::CopyCL0c2Gm(const TensorC& tens | |||
| 364 | auto layoutL0C = AscendC::Te::MakeFrameLayout<AscendC::Te::NZLayoutPtn, AscendC::Std::Int<L0C_C0>>(param.mL1Size, param.nL1Size); | 360 | auto layoutL0C = AscendC::Te::MakeFrameLayout<AscendC::Te::NZLayoutPtn, AscendC::Std::Int<L0C_C0>>(param.mL1Size, param.nL1Size); |
| 365 | auto tensorL0C = AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0C, float>(0), layoutL0C); | 361 | auto tensorL0C = AscendC::Te::MakeTensor(AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0C, float>(0), layoutL0C); |
| 366 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 362 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 367 | - AscendC::Te::Copy(CopyL0C2GM, tensorC, tensorL0C, FP32_64_AS_UINT64, AscendC::Te::FixpipeParams{/*unitflag*/ 3}); | 363 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{/*unitflag*/ 3}), tensorC, tensorL0C, FP32_64_AS_UINT64); |
| 368 | } | 364 | } |
| 369 | 365 | ||
| 370 | WQBMM_CUBE_COMPUTE_TEMPLATE_PARAM | 366 | WQBMM_CUBE_COMPUTE_TEMPLATE_PARAM |
MSamples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/include/kernel/quant_grouped_matmul_mx_kernel.h+1-1
| @@ -153,7 +153,7 @@ private: | |||
| 153 | constexpr uint32_t UB_ALIGN_SIZE = 32; | 153 | constexpr uint32_t UB_ALIGN_SIZE = 32; |
| 154 | constexpr uint32_t MAX_REPEAT_TIMES = 255; | 154 | constexpr uint32_t MAX_REPEAT_TIMES = 255; |
| 155 | constexpr uint32_t ONE_BLK_SIZE = 32; | 155 | constexpr uint32_t ONE_BLK_SIZE = 32; |
| 156 | - uint64_t usedBlockNum = AscendC::GetBlockNum() * AscendC::GetTaskRation(); | 156 | + uint64_t usedBlockNum = AscendC::GetBlockNum() * AscendC::GetTaskRatio(); |
| 157 | 157 | ||
| 158 | // Calculate init buffer size in elements | 158 | // Calculate init buffer size in elements |
| 159 | uint64_t ySize = static_cast<uint64_t>(m) * n; | 159 | uint64_t ySize = static_cast<uint64_t>(m) * n; |
MSamples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/include/tile/copy_scale_l1_to_l0b.h+0-11
| @@ -46,14 +46,3 @@ struct AscendC::Te::CopyTraits<::Tile::CopyL12L0MxScaleB3510> | |||
| 46 | ::Tile::CopyL12L0MxScaleB3510, CopyL12L0BTraitDefault, ::Tile::CopyL12L0MxScaleB3510, | 46 | ::Tile::CopyL12L0MxScaleB3510, CopyL12L0BTraitDefault, ::Tile::CopyL12L0MxScaleB3510, |
| 47 | CopyL12L0BTraitDefault> {}; | 47 | CopyL12L0BTraitDefault> {}; |
| 48 | 48 | ||
| 49 | -namespace AscendC::Te { | ||
| 50 | - | ||
| 51 | -constexpr LoadDataTrait LOAD_DATA_B_TRAIT{true}; | ||
| 52 | - | ||
| 53 | -struct LoadData2BTrait { | ||
| 54 | - using TraitType = LoadDataTrait; | ||
| 55 | - static constexpr const TraitType value = LOAD_DATA_B_TRAIT; | ||
| 56 | -}; | ||
| 57 | - | ||
| 58 | -} // namespace AscendC::Te | ||
| 59 | - | ||
| @@ -203,8 +203,9 @@ public: | |||
| 203 | AscendC::Te::MmadParams mmadParams(curML1, curNL1, curK0, unitFlag, cmatrixInitVal); | 203 | AscendC::Te::MmadParams mmadParams(curML1, curNL1, curK0, unitFlag, cmatrixInitVal); |
| 204 | // Pass custom Trait type in mmad | 204 | // Pass custom Trait type in mmad |
| 205 | AscendC::Te::Mmad( | 205 | AscendC::Te::Mmad( |
| 206 | - AscendC::Te::MmadAtom<AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, MmadMmTraitConfig>>{}, | 206 | + AscendC::Te::MmadAtom<AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, MmadMmTraitConfig>>{} |
| 207 | - tensorL0C, tensorAL0, tensorBL0, mmadParams); | 207 | + .with(mmadParams), |
| 208 | + tensorL0C, tensorAL0, tensorBL0); | ||
| 208 | 209 | ||
| 209 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0PingPong_ & 0x1); | 210 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0PingPong_ & 0x1); |
| 210 | l0PingPong_++; | 211 | l0PingPong_++; |
| @@ -214,9 +215,9 @@ public: | |||
| 214 | // Depending on checkIsSkScene, decide to move out to GM or WorkSpace | 215 | // Depending on checkIsSkScene, decide to move out to GM or WorkSpace |
| 215 | if (checkIsSkScene) { | 216 | if (checkIsSkScene) { |
| 216 | AscendC::Te::Copy( | 217 | AscendC::Te::Copy( |
| 217 | - CopyL0C2GM, gmWorkSpace, tensorL0C, AscendC::Te::FixpipeParams(FINAL_ACCUMULATION)); | 218 | + CopyL0C2GM.with(AscendC::Te::FixpipeParams(FINAL_ACCUMULATION)), gmWorkSpace, tensorL0C); |
| 218 | } else { | 219 | } else { |
| 219 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams(FINAL_ACCUMULATION)); | 220 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams(FINAL_ACCUMULATION)), gmC, tensorL0C); |
| 220 | } | 221 | } |
| 221 | } | 222 | } |
| 222 | AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(l1BufId); | 223 | AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(l1BufId); |
| @@ -219,8 +219,9 @@ public: | |||
| 219 | 219 | ||
| 220 | // Pass custom Trait type in mmad | 220 | // Pass custom Trait type in mmad |
| 221 | AscendC::Te::Mmad( | 221 | AscendC::Te::Mmad( |
| 222 | - AscendC::Te::MmadAtom<AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, MmadMmTraitConfig>>{}, | 222 | + AscendC::Te::MmadAtom<AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, MmadMmTraitConfig>>{} |
| 223 | - tensorL0C, tensorAL0, tensorBL0, mmadParams); | 223 | + .with(mmadParams), |
| 224 | + tensorL0C, tensorAL0, tensorBL0); | ||
| 224 | 225 | ||
| 225 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(static_cast<uint16_t>(mte1Flag)); | 226 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(static_cast<uint16_t>(mte1Flag)); |
| 226 | l0PingPong_++; | 227 | l0PingPong_++; |
| @@ -237,7 +238,7 @@ public: | |||
| 237 | AscendC::Te::FixpipeParams fixpParams; | 238 | AscendC::Te::FixpipeParams fixpParams; |
| 238 | fixpParams.unitFlag = enableL0cPingPong_ ? 0 : FINAL_ACCUMULATION; | 239 | fixpParams.unitFlag = enableL0cPingPong_ ? 0 : FINAL_ACCUMULATION; |
| 239 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 240 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 240 | - AscendC::Te::Copy(copyL0C2GM, gmC, tensorL0C, fixpParams); | 241 | + AscendC::Te::Copy(copyL0C2GM.with(fixpParams), gmC, tensorL0C); |
| 241 | 242 | ||
| 242 | if (enableL0cPingPong_) { | 243 | if (enableL0cPingPong_) { |
| 243 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(l0cPingPong_ & 0x1); | 244 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(l0cPingPong_ & 0x1); |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_hifp8_block_mmad_swat.h+7-6
| @@ -202,11 +202,11 @@ public: | |||
| 202 | 202 | ||
| 203 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 203 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 204 | if constexpr (isScalarScale) { | 204 | if constexpr (isScalarScale) { |
| 205 | - AscendC::Te::Copy(copyL0C2GM, gmC, c1Local, scalarScale_, | 205 | + AscendC::Te::Copy(copyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, c1Local, |
| 206 | - AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 206 | + scalarScale_); |
| 207 | } else { | 207 | } else { |
| 208 | - AscendC::Te::Copy(copyL0C2GM, gmC, c1Local, tensorX2L1, | 208 | + AscendC::Te::Copy(copyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, c1Local, |
| 209 | - AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 209 | + tensorX2L1); |
| 210 | } | 210 | } |
| 211 | 211 | ||
| 212 | if (enableL0cPingPong_) { | 212 | if (enableL0cPingPong_) { |
| @@ -270,8 +270,9 @@ private: | |||
| 270 | } | 270 | } |
| 271 | AscendC::Te::Mmad( | 271 | AscendC::Te::Mmad( |
| 272 | AscendC::Te::MmadAtom< | 272 | AscendC::Te::MmadAtom< |
| 273 | - AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, AscendC::Te::MmadTraitDefault>>{}, | 273 | + AscendC::Te::MmadTraits<AscendC::Te::MmadOperation, AscendC::Te::MmadTraitDefault>>{} |
| 274 | - c1Local, l0aLocal, l0bLocal, mmadParams); | 274 | + .with(mmadParams), |
| 275 | + c1Local, l0aLocal, l0bLocal); | ||
| 275 | 276 | ||
| 276 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0PingPong_ & 0x1); | 277 | AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(l0PingPong_ & 0x1); |
| 277 | l0PingPong_++; | 278 | l0PingPong_++; |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_hifp8_block_scheduler_swat.h+1-1
| @@ -40,7 +40,7 @@ template <class ProblemShape_, bool TransA_, bool TransB_> | |||
| 40 | int64_t nBaseTailLast_{0}; | 40 | int64_t nBaseTailLast_{0}; |
| 41 | int64_t mCoreNum_{0}; | 41 | int64_t mCoreNum_{0}; |
| 42 | int64_t mTailCoreNum_{0}; | 42 | int64_t mTailCoreNum_{0}; |
| 43 | - int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRation()}; | 43 | + int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRatio()}; |
| 44 | int64_t blockNum_{AscendC::GetBlockNum()}; | 44 | int64_t blockNum_{AscendC::GetBlockNum()}; |
| 45 | int64_t startBlockIdx_{0}; | 45 | int64_t startBlockIdx_{0}; |
| 46 | int64_t endBlockIdx_{0}; | 46 | int64_t endBlockIdx_{0}; |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_mx_block_mmad_a_full_load.h+3-5
| @@ -325,8 +325,7 @@ public: | |||
| 325 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( | 325 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( |
| 326 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); | 326 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); |
| 327 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 327 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 328 | - AscendC::Te::Copy( | 328 | + CopyL12L0MxScaleA3510.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 329 | - CopyL12L0MxScaleA3510, tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 330 | 329 | ||
| 331 | auto layoutBL0 = | 330 | auto layoutBL0 = |
| 332 | AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); | 331 | AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); |
| @@ -342,8 +341,7 @@ public: | |||
| 342 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( | 341 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( |
| 343 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); | 342 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); |
| 344 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 343 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 345 | - AscendC::Te::Copy( | 344 | + CopyL12L0MxScaleB3510.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 346 | - CopyL12L0MxScaleB3510, tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 347 | 345 | ||
| 348 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 346 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 349 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 347 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| @@ -380,7 +378,7 @@ public: | |||
| 380 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 378 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 381 | // The whole block accumulates into one L0C tile, which is flushed once | 379 | // The whole block accumulates into one L0C tile, which is flushed once |
| 382 | // after all K chunks have contributed. | 380 | // after all K chunks have contributed. |
| 383 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 381 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 384 | if (enableL0cPingPong_) { | 382 | if (enableL0cPingPong_) { |
| 385 | l0cPingPong_++; | 383 | l0cPingPong_++; |
| 386 | } | 384 | } |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_mx_block_mmad_swat.h+3-5
| @@ -311,8 +311,7 @@ public: | |||
| 311 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice( | 311 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice( |
| 312 | AscendC::Te::MakeCoord(0, coordScaleKL1), AscendC::Te::MakeShape(curM, kL1ScaleSize_)); | 312 | AscendC::Te::MakeCoord(0, coordScaleKL1), AscendC::Te::MakeShape(curM, kL1ScaleSize_)); |
| 313 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 313 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 314 | - AscendC::Te::Copy( | 314 | + CopyL12L0MxScaleA3510.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 315 | - CopyL12L0MxScaleA3510, tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 316 | 315 | ||
| 317 | auto layoutScaleBL0 = | 316 | auto layoutScaleBL0 = |
| 318 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 317 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| @@ -328,8 +327,7 @@ public: | |||
| 328 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice( | 327 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice( |
| 329 | AscendC::Te::MakeCoord(coordScaleKL1, 0), AscendC::Te::MakeShape(kL1ScaleSize_, curN)); | 328 | AscendC::Te::MakeCoord(coordScaleKL1, 0), AscendC::Te::MakeShape(kL1ScaleSize_, curN)); |
| 330 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 329 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 331 | - AscendC::Te::Copy( | 330 | + CopyL12L0MxScaleB3510.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 332 | - CopyL12L0MxScaleB3510, tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 333 | 331 | ||
| 334 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 332 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 335 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 333 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| @@ -367,7 +365,7 @@ public: | |||
| 367 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 365 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 368 | // The whole block accumulates into one L0C tile, which is flushed once | 366 | // The whole block accumulates into one L0C tile, which is flushed once |
| 369 | // after all K chunks have contributed. | 367 | // after all K chunks have contributed. |
| 370 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 368 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 371 | if (enableL0cPingPong_) { | 369 | if (enableL0cPingPong_) { |
| 372 | l0cPingPong_++; | 370 | l0cPingPong_++; |
| 373 | } | 371 | } |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_mx_block_mmad_swat_4_buffer.h+3-5
| @@ -311,8 +311,7 @@ public: | |||
| 311 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice( | 311 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice( |
| 312 | AscendC::Te::MakeCoord(0, coordScaleKL1), AscendC::Te::MakeShape(curM, kL1ScaleSize_)); | 312 | AscendC::Te::MakeCoord(0, coordScaleKL1), AscendC::Te::MakeShape(curM, kL1ScaleSize_)); |
| 313 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 313 | auto CopyL12L0MxScaleA3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 314 | - AscendC::Te::Copy( | 314 | + CopyL12L0MxScaleA3510.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 315 | - CopyL12L0MxScaleA3510, tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 316 | 315 | ||
| 317 | auto layoutScaleBL0 = | 316 | auto layoutScaleBL0 = |
| 318 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 317 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| @@ -328,8 +327,7 @@ public: | |||
| 328 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice( | 327 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice( |
| 329 | AscendC::Te::MakeCoord(coordScaleKL1, 0), AscendC::Te::MakeShape(kL1ScaleSize_, curN)); | 328 | AscendC::Te::MakeCoord(coordScaleKL1, 0), AscendC::Te::MakeShape(kL1ScaleSize_, curN)); |
| 330 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 329 | auto CopyL12L0MxScaleB3510 = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 331 | - AscendC::Te::Copy( | 330 | + CopyL12L0MxScaleB3510.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 332 | - CopyL12L0MxScaleB3510, tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 333 | 331 | ||
| 334 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 332 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| 335 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); | 333 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0BufId); |
| @@ -367,7 +365,7 @@ public: | |||
| 367 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 365 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 368 | // The whole block accumulates into one L0C tile, which is flushed once | 366 | // The whole block accumulates into one L0C tile, which is flushed once |
| 369 | // after all K chunks have contributed. | 367 | // after all K chunks have contributed. |
| 370 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 368 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 371 | if (enableL0cPingPong_) { | 369 | if (enableL0cPingPong_) { |
| 372 | l0cPingPong_++; | 370 | l0cPingPong_++; |
| 373 | } | 371 | } |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_mx_block_scheduler_a_full_load.h+1-1
| @@ -38,7 +38,7 @@ public: | |||
| 38 | int64_t nBaseTailMain_{0}; | 38 | int64_t nBaseTailMain_{0}; |
| 39 | int64_t mBaseTailLast_{0}; | 39 | int64_t mBaseTailLast_{0}; |
| 40 | int64_t nBaseTailLast_{0}; | 40 | int64_t nBaseTailLast_{0}; |
| 41 | - int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRation()}; | 41 | + int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRatio()}; |
| 42 | int64_t blockNum_{AscendC::GetBlockNum()}; | 42 | int64_t blockNum_{AscendC::GetBlockNum()}; |
| 43 | int64_t endBlockIdx_{0}; | 43 | int64_t endBlockIdx_{0}; |
| 44 | int64_t roundIdx_{0}; | 44 | int64_t roundIdx_{0}; |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_mx_block_scheduler_swat.h+1-1
| @@ -40,7 +40,7 @@ public: | |||
| 40 | int64_t nBaseTailLast_{0}; | 40 | int64_t nBaseTailLast_{0}; |
| 41 | int64_t mCoreNum_{0}; | 41 | int64_t mCoreNum_{0}; |
| 42 | int64_t mTailCoreNum_{0}; | 42 | int64_t mTailCoreNum_{0}; |
| 43 | - int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRation()}; | 43 | + int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRatio()}; |
| 44 | int64_t blockNum_{AscendC::GetBlockNum()}; | 44 | int64_t blockNum_{AscendC::GetBlockNum()}; |
| 45 | int64_t startBlockIdx_{0}; | 45 | int64_t startBlockIdx_{0}; |
| 46 | int64_t endBlockIdx_{0}; | 46 | int64_t endBlockIdx_{0}; |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/quant_matmul_mx_block_scheduler_swat_4_buffer.h+1-1
| @@ -41,7 +41,7 @@ public: | |||
| 41 | int64_t nBaseTailLast_{0}; | 41 | int64_t nBaseTailLast_{0}; |
| 42 | int64_t mCoreNum_{0}; | 42 | int64_t mCoreNum_{0}; |
| 43 | int64_t mTailCoreNum_{0}; | 43 | int64_t mTailCoreNum_{0}; |
| 44 | - int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRation()}; | 44 | + int64_t blockIdx_{AscendC::GetBlockIdx() / AscendC::GetTaskRatio()}; |
| 45 | int64_t blockNum_{AscendC::GetBlockNum()}; | 45 | int64_t blockNum_{AscendC::GetBlockNum()}; |
| 46 | int64_t startBlockIdx_{0}; | 46 | int64_t startBlockIdx_{0}; |
| 47 | int64_t endBlockIdx_{0}; | 47 | int64_t endBlockIdx_{0}; |
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/weight_quant_matmul_mxfp8fp4_block_mmad_swat.h+4-7
| @@ -385,9 +385,7 @@ private: | |||
| 385 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( | 385 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( |
| 386 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, ScaleAType>(l0Offset), layoutScaleAL0); | 386 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, ScaleAType>(l0Offset), layoutScaleAL0); |
| 387 | auto copyScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 387 | auto copyScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 388 | - AscendC::Te::Copy( | 388 | + copyScaleA.Call(tensorScaleAL0, tensorScaleAL1, AscendC::Te::MakeCoord(0, static_cast<int64_t>(kOffsetInScaleWindow))); |
| 389 | - copyScaleA, tensorScaleAL0, tensorScaleAL1, | ||
| 390 | - AscendC::Te::MakeCoord(0, static_cast<int64_t>(kOffsetInScaleWindow))); | ||
| 391 | 389 | ||
| 392 | auto layoutScaleBL0 = | 390 | auto layoutScaleBL0 = |
| 393 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<MXFP_MULTI_BASE_SIZE>>( | 391 | AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<MXFP_MULTI_BASE_SIZE>>( |
| @@ -395,9 +393,7 @@ private: | |||
| 395 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( | 393 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( |
| 396 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, ScaleBType>(l0Offset), layoutScaleBL0); | 394 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, ScaleBType>(l0Offset), layoutScaleBL0); |
| 397 | auto copyScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 395 | auto copyScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 398 | - AscendC::Te::Copy( | 396 | + copyScaleB.Call(tensorScaleBL0, tensorScaleBL1, AscendC::Te::MakeCoord(static_cast<int64_t>(kOffsetInScaleWindow), 0)); |
| 399 | - copyScaleB, tensorScaleBL0, tensorScaleBL1, | ||
| 400 | - AscendC::Te::MakeCoord(static_cast<int64_t>(kOffsetInScaleWindow), 0)); | ||
| 401 | } | 397 | } |
| 402 | 398 | ||
| 403 | template <typename TensorC, typename TensorL0C> | 399 | template <typename TensorC, typename TensorL0C> |
| @@ -406,7 +402,8 @@ private: | |||
| 406 | constexpr uint64_t FP32_64_AS_UINT64 = 0x42800000; | 402 | constexpr uint64_t FP32_64_AS_UINT64 = 0x42800000; |
| 407 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 403 | auto copyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 408 | AscendC::Te::Copy( | 404 | AscendC::Te::Copy( |
| 409 | - copyL0C2GM, tensorC, tensorL0C, FP32_64_AS_UINT64, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 405 | + copyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), tensorC, tensorL0C, |
| 406 | + FP32_64_AS_UINT64); | ||
| 410 | } | 407 | } |
| 411 | 408 | ||
| 412 | __aicore__ inline void PostProcess( | 409 | __aicore__ inline void PostProcess( |
MSamples/2_Performance/matmul_story/matmul_recipes/include/epilogue/matmul_block_epilogue_streamk.h+4-4
| @@ -102,7 +102,7 @@ public: | |||
| 102 | kCnt_ = AscendC::Te::Get<MNK_K>(coordInAiv); | 102 | kCnt_ = AscendC::Te::Get<MNK_K>(coordInAiv); |
| 103 | usedCoreNum_ = usedCoreNum; | 103 | usedCoreNum_ = usedCoreNum; |
| 104 | // Decrease tile size of per vector core to prevent data race of cube and vector | 104 | // Decrease tile size of per vector core to prevent data race of cube and vector |
| 105 | - aivMte2Num_ = checkIsSkScene ? AscendC::GetTaskRation() : AscendC::BLOCK_CUBE; | 105 | + aivMte2Num_ = checkIsSkScene ? AscendC::GetTaskRatio() : AscendC::BLOCK_CUBE; |
| 106 | cGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ OutType*>(params.cGmAddr)); | 106 | cGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ OutType*>(params.cGmAddr)); |
| 107 | workspaceGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ WorkspaceType*>(params.workspaceGmAddr)); | 107 | workspaceGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ WorkspaceType*>(params.workspaceGmAddr)); |
| 108 | ICachePreLoad(NUM_TWO); | 108 | ICachePreLoad(NUM_TWO); |
| @@ -115,8 +115,8 @@ public: | |||
| 115 | 115 | ||
| 116 | __aicore__ inline void UpdateAivBasicIndex() | 116 | __aicore__ inline void UpdateAivBasicIndex() |
| 117 | { | 117 | { |
| 118 | - uint64_t newBlockIdx = AscendC::GetBlockIdx() / (AscendC::GetTaskRation() * kCnt_); | 118 | + uint64_t newBlockIdx = AscendC::GetBlockIdx() / (AscendC::GetTaskRatio() * kCnt_); |
| 119 | - aivParams_.kCntIndex = AscendC::GetBlockIdx() % (AscendC::GetTaskRation() * kCnt_); | 119 | + aivParams_.kCntIndex = AscendC::GetBlockIdx() % (AscendC::GetTaskRatio() * kCnt_); |
| 120 | 120 | ||
| 121 | aivParams_.indexParams = newBlockIdx; | 121 | aivParams_.indexParams = newBlockIdx; |
| 122 | uint64_t cGmIndex = aivParams_.indexParams + (mCnt_ * nCnt_ - mCnt_ * nCnt_ % usedCoreNum_); | 122 | uint64_t cGmIndex = aivParams_.indexParams + (mCnt_ * nCnt_ - mCnt_ * nCnt_ % usedCoreNum_); |
| @@ -151,7 +151,7 @@ public: | |||
| 151 | __aicore__ inline void UpdateAivParams(uint64_t index) | 151 | __aicore__ inline void UpdateAivParams(uint64_t index) |
| 152 | { | 152 | { |
| 153 | mBurstBase_ = CeilAlign( | 153 | mBurstBase_ = CeilAlign( |
| 154 | - CeilDiv(aivParams_.curML1InAiv, kCnt_ * AscendC::GetTaskRation()), | 154 | + CeilDiv(aivParams_.curML1InAiv, kCnt_ * AscendC::GetTaskRatio()), |
| 155 | CeilDiv(BLOCK_BYTE_SIZE, aivParams_.curAlignedNInAiv)); | 155 | CeilDiv(BLOCK_BYTE_SIZE, aivParams_.curAlignedNInAiv)); |
| 156 | uint64_t mBurstCnt = CeilDiv(aivParams_.curML1InAiv, mBurstBase_); | 156 | uint64_t mBurstCnt = CeilDiv(aivParams_.curML1InAiv, mBurstBase_); |
| 157 | uint64_t mBurstTail = aivParams_.curML1InAiv - (mBurstCnt - 1) * mBurstBase_; | 157 | uint64_t mBurstTail = aivParams_.curML1InAiv - (mBurstCnt - 1) * mBurstBase_; |
| @@ -15,7 +15,7 @@ | |||
| 15 | 15 | ||
| 16 | 16 | ||
| 17 | 17 | ||
| 18 | -#include "blaze/utils/common_utils.h" | 18 | +#include "blaze/gemm/utils/common_utils.h" |
| 19 | 19 | ||
| 20 | namespace Kernel { | 20 | namespace Kernel { |
| 21 | 21 | ||
| @@ -204,7 +204,7 @@ __aicore__ inline void MatmulKernelStreamK<ProblemShape, BlockMmad, BlockSchedul | |||
| 204 | if ASCEND_IS_AIV { | 204 | if ASCEND_IS_AIV { |
| 205 | uint64_t lastLoopTotalCnt = (mTileNum * nTileNum % usedCoreNum_) * skKTileNum; | 205 | uint64_t lastLoopTotalCnt = (mTileNum * nTileNum % usedCoreNum_) * skKTileNum; |
| 206 | uint64_t curBlockIdxInAiv = AscendC::GetBlockIdx(); | 206 | uint64_t curBlockIdxInAiv = AscendC::GetBlockIdx(); |
| 207 | - if (curBlockIdxInAiv >= lastLoopTotalCnt * AscendC::GetTaskRation()) { | 207 | + if (curBlockIdxInAiv >= lastLoopTotalCnt * AscendC::GetTaskRatio()) { |
| 208 | AscendC::CrossCoreWaitFlag<AIC_SYNC_AIV_MODE_4, PIPE_MTE3>(AIC_SYNC_AIV_FLAG); | 208 | AscendC::CrossCoreWaitFlag<AIC_SYNC_AIV_MODE_4, PIPE_MTE3>(AIC_SYNC_AIV_FLAG); |
| 209 | AscendC::SyncAll(); | 209 | AscendC::SyncAll(); |
| 210 | return; | 210 | return; |
MSamples/2_Performance/matmul_story/matmul_recipes/include/kernel/weight_quant_matmul_mxfp8fp4_kernel_swat.h+1-1
| @@ -333,7 +333,7 @@ private: | |||
| 333 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::GM>(reinterpret_cast<__gm__ BType*>(params.mmad.bGmAddr)), | 333 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::GM>(reinterpret_cast<__gm__ BType*>(params.mmad.bGmAddr)), |
| 334 | LayoutB{}(static_cast<int64_t>(params.problemShape.k), static_cast<int64_t>(params.problemShape.n))); | 334 | LayoutB{}(static_cast<int64_t>(params.problemShape.k), static_cast<int64_t>(params.problemShape.n))); |
| 335 | uint64_t tileNum = scheduler.GetTileNum(); | 335 | uint64_t tileNum = scheduler.GetTileNum(); |
| 336 | - uint64_t curBlockIdx = AscendC::GetBlockIdx() / AscendC::GetTaskRation(); | 336 | + uint64_t curBlockIdx = AscendC::GetBlockIdx() / AscendC::GetTaskRatio(); |
| 337 | typename BlockPrologue::Params prologueParams{ | 337 | typename BlockPrologue::Params prologueParams{ |
| 338 | params.mmad.bGmAddr, static_cast<uint64_t>(AscendC::Std::get<1>(params.mmad.l1TileShape)), | 338 | params.mmad.bGmAddr, static_cast<uint64_t>(AscendC::Std::get<1>(params.mmad.l1TileShape)), |
| 339 | static_cast<uint64_t>(AscendC::Std::get<2>(params.mmad.l1TileShape)), params.mmad.kBubSize, | 339 | static_cast<uint64_t>(AscendC::Std::get<2>(params.mmad.l1TileShape)), params.mmad.kBubSize, |
| @@ -175,8 +175,7 @@ public: | |||
| 175 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, 0), | 175 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, 0), |
| 176 | AscendC::Te::MakeShape(curM, CeilDiv(curKL1, MXFP_DIVISOR_SIZE) * MXFP_MULTI_BASE_SIZE)); | 176 | AscendC::Te::MakeShape(curM, CeilDiv(curKL1, MXFP_DIVISOR_SIZE) * MXFP_MULTI_BASE_SIZE)); |
| 177 | auto copyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 177 | auto copyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 178 | - AscendC::Te::Copy( | 178 | + copyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 179 | - copyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 180 | 179 | ||
| 181 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 180 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| 182 | CeilDiv(curKL0, MXFP_DIVISOR_SIZE) * MXFP_MULTI_BASE_SIZE, curN); | 181 | CeilDiv(curKL0, MXFP_DIVISOR_SIZE) * MXFP_MULTI_BASE_SIZE, curN); |
| @@ -185,8 +184,7 @@ public: | |||
| 185 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(0, 0), | 184 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(0, 0), |
| 186 | AscendC::Te::MakeShape(CeilDiv(curKL1, MXFP_DIVISOR_SIZE) * MXFP_MULTI_BASE_SIZE, curN)); | 185 | AscendC::Te::MakeShape(CeilDiv(curKL1, MXFP_DIVISOR_SIZE) * MXFP_MULTI_BASE_SIZE, curN)); |
| 187 | auto copyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 186 | auto copyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 188 | - AscendC::Te::Copy( | 187 | + copyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 189 | - copyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 190 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(0); | 188 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(0); |
| 191 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(0); | 189 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(0); |
| 192 | 190 | ||
MSamples/2_Performance/matmul_story/matmul_tutorials/1_pingpong/include/block/block_mmad_mx_base.h+3-7
| @@ -225,9 +225,7 @@ public: | |||
| 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), | 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), |
| 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); | 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); |
| 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 228 | - AscendC::Te::Copy( | 228 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 229 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 230 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 231 | 229 | ||
| 232 | // ScaleB: L1 -> L0B | 230 | // ScaleB: L1 -> L0B |
| 233 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 231 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| @@ -241,9 +239,7 @@ public: | |||
| 241 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), | 239 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), |
| 242 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); | 240 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); |
| 243 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 241 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 244 | - AscendC::Te::Copy( | 242 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 245 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 246 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 247 | 243 | ||
| 248 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 244 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 249 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 245 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -280,7 +276,7 @@ AscendC::Te::Mmad( | |||
| 280 | 276 | ||
| 281 | // L0C -> GM (unitFlag = 0) | 277 | // L0C -> GM (unitFlag = 0) |
| 282 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 278 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 283 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{0}); | 279 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{0}), gmC, tensorL0C); |
| 284 | 280 | ||
| 285 | // Wait for fixpipe to finish before next Cube iteration | 281 | // Wait for fixpipe to finish before next Cube iteration |
| 286 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(0); | 282 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(0); |
MSamples/2_Performance/matmul_story/matmul_tutorials/2_block_swat/include/block/block_mmad_mx_swat.h+3-7
| @@ -225,9 +225,7 @@ public: | |||
| 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), | 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), |
| 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); | 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); |
| 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 228 | - AscendC::Te::Copy( | 228 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 229 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 230 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 231 | 229 | ||
| 232 | // ScaleB: L1 -> L0B | 230 | // ScaleB: L1 -> L0B |
| 233 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 231 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| @@ -241,9 +239,7 @@ public: | |||
| 241 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), | 239 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), |
| 242 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); | 240 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); |
| 243 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 241 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 244 | - AscendC::Te::Copy( | 242 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 245 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 246 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 247 | 243 | ||
| 248 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 244 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 249 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 245 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -280,7 +276,7 @@ AscendC::Te::Mmad( | |||
| 280 | 276 | ||
| 281 | // L0C -> GM (unitFlag = 0) | 277 | // L0C -> GM (unitFlag = 0) |
| 282 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 278 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 283 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{0}); | 279 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{0}), gmC, tensorL0C); |
| 284 | 280 | ||
| 285 | // Wait for fixpipe to finish before next Cube iteration | 281 | // Wait for fixpipe to finish before next Cube iteration |
| 286 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(0); | 282 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(0); |
MSamples/2_Performance/matmul_story/matmul_tutorials/3_last_round_tile_balance/include/block/block_mmad_mx_swat.h+3-7
| @@ -225,9 +225,7 @@ public: | |||
| 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), | 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), |
| 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); | 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); |
| 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 228 | - AscendC::Te::Copy( | 228 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 229 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 230 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 231 | 229 | ||
| 232 | // ScaleB: L1 -> L0B | 230 | // ScaleB: L1 -> L0B |
| 233 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 231 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| @@ -241,9 +239,7 @@ public: | |||
| 241 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), | 239 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), |
| 242 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); | 240 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); |
| 243 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 241 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 244 | - AscendC::Te::Copy( | 242 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 245 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 246 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 247 | 243 | ||
| 248 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 244 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 249 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 245 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -280,7 +276,7 @@ AscendC::Te::Mmad( | |||
| 280 | 276 | ||
| 281 | // L0C -> GM (unitFlag = 0) | 277 | // L0C -> GM (unitFlag = 0) |
| 282 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 278 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 283 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{0}); | 279 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{0}), gmC, tensorL0C); |
| 284 | 280 | ||
| 285 | // Wait for fixpipe to finish before next Cube iteration | 281 | // Wait for fixpipe to finish before next Cube iteration |
| 286 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(0); | 282 | AscendC::SetFlag<AscendC::HardEvent::FIX_M>(0); |
MSamples/2_Performance/matmul_story/matmul_tutorials/4_unit_flag/include/block/block_mmad_mx_swat.h+3-7
| @@ -225,9 +225,7 @@ public: | |||
| 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), | 225 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), |
| 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); | 226 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); |
| 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 227 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 228 | - AscendC::Te::Copy( | 228 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 229 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 230 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 231 | 229 | ||
| 232 | // ScaleB: L1 -> L0B | 230 | // ScaleB: L1 -> L0B |
| 233 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 231 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| @@ -241,9 +239,7 @@ public: | |||
| 241 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), | 239 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), |
| 242 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); | 240 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); |
| 243 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 241 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 244 | - AscendC::Te::Copy( | 242 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 245 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 246 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 247 | 243 | ||
| 248 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 244 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 249 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 245 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -278,7 +274,7 @@ AscendC::Te::Mmad( | |||
| 278 | 274 | ||
| 279 | // L0C -> GM | 275 | // L0C -> GM |
| 280 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 276 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 281 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 277 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 282 | 278 | ||
| 283 | if (enableL0cPingPong_) { | 279 | if (enableL0cPingPong_) { |
| 284 | l0cPingPong_++; | 280 | l0cPingPong_++; |
MSamples/2_Performance/matmul_story/matmul_tutorials/5_halfl1_ping_halfl1_pong/include/block/block_mmad_mx_base.h+3-7
| @@ -224,9 +224,7 @@ public: | |||
| 224 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), | 224 | auto tensorBlockScaleAL1 = tensorScaleAL1.Slice(AscendC::Te::MakeCoord(0, coordScaleKL1), |
| 225 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); | 225 | AscendC::Te::MakeShape(curM, CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL)); |
| 226 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 226 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 227 | - AscendC::Te::Copy( | 227 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 228 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 229 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 230 | 228 | ||
| 231 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( | 229 | auto layoutScaleBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::NNLayoutPtn, AscendC::Std::Int<SCALE_C0>>( |
| 232 | CeilDiv(curKL0, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN); | 230 | CeilDiv(curKL0, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN); |
| @@ -239,9 +237,7 @@ public: | |||
| 239 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), | 237 | auto tensorBlockScaleBL1 = tensorScaleBL1.Slice(AscendC::Te::MakeCoord(coordScaleKL1, 0), |
| 240 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); | 238 | AscendC::Te::MakeShape(CeilDiv(kL1_, MXFP_DIVISOR_SIZE_LOCAL) * MXFP_MULTI_BASE_SIZE_LOCAL, curN)); |
| 241 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 239 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 242 | - AscendC::Te::Copy( | 240 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 243 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 244 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 245 | 241 | ||
| 246 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 242 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 247 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 243 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -275,7 +271,7 @@ AscendC::Te::Mmad( | |||
| 275 | } | 271 | } |
| 276 | 272 | ||
| 277 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 273 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 278 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 274 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 279 | 275 | ||
| 280 | if (enableL0cPingPong_) { | 276 | if (enableL0cPingPong_) { |
| 281 | l0cPingPong_++; | 277 | l0cPingPong_++; |
MSamples/2_Performance/matmul_story/matmul_tutorials/6_scale_memory_access_coalescing/include/block/block_mmad_mx_base.h+3-7
| @@ -152,9 +152,7 @@ public: | |||
| 152 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( | 152 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( |
| 153 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); | 153 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); |
| 154 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 154 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 155 | - AscendC::Te::Copy( | 155 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, scaleKL1IterOffset + kL0Offset)); |
| 156 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 157 | - AscendC::Te::MakeCoord(0, scaleKL1IterOffset + kL0Offset)); | ||
| 158 | 156 | ||
| 159 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); | 157 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); |
| 160 | auto tensorBL0 = | 158 | auto tensorBL0 = |
| @@ -166,9 +164,7 @@ public: | |||
| 166 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( | 164 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( |
| 167 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); | 165 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); |
| 168 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 166 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 169 | - AscendC::Te::Copy( | 167 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(scaleKL1IterOffset + kL0Offset, 0)); |
| 170 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 171 | - AscendC::Te::MakeCoord(scaleKL1IterOffset + kL0Offset, 0)); | ||
| 172 | 168 | ||
| 173 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 169 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 174 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 170 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -288,7 +284,7 @@ AscendC::Te::Mmad( | |||
| 288 | } | 284 | } |
| 289 | 285 | ||
| 290 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 286 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 291 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 287 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 292 | if (enableL0cPingPong_) { | 288 | if (enableL0cPingPong_) { |
| 293 | l0cPingPong_++; | 289 | l0cPingPong_++; |
| 294 | } | 290 | } |
MSamples/2_Performance/matmul_story/matmul_tutorials/7_fullload/include/block/quant_matmul_mxfp4_block_mmad_a_full_load.h+3-7
| @@ -151,9 +151,7 @@ public: | |||
| 151 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( | 151 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( |
| 152 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); | 152 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); |
| 153 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 153 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 154 | - AscendC::Te::Copy( | 154 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 155 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 156 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 157 | 155 | ||
| 158 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); | 156 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); |
| 159 | auto tensorBL0 = | 157 | auto tensorBL0 = |
| @@ -165,9 +163,7 @@ public: | |||
| 165 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( | 163 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( |
| 166 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); | 164 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); |
| 167 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 165 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 168 | - AscendC::Te::Copy( | 166 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 169 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 170 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 171 | 167 | ||
| 172 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 168 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 173 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 169 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -286,7 +282,7 @@ AscendC::Te::Mmad( | |||
| 286 | } | 282 | } |
| 287 | 283 | ||
| 288 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 284 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 289 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 285 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 290 | if (enableL0cPingPong_) { | 286 | if (enableL0cPingPong_) { |
| 291 | l0cPingPong_++; | 287 | l0cPingPong_++; |
| 292 | } | 288 | } |
MSamples/2_Performance/matmul_story/matmul_tutorials/7_fullload/include/block/quant_matmul_mxfp4_block_mmad_swat.h+3-7
| @@ -150,9 +150,7 @@ public: | |||
| 150 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( | 150 | auto tensorScaleAL0 = AscendC::Te::MakeTensor( |
| 151 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); | 151 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0A, fp8_e8m0_t>(l0Offset), layoutScaleAL0); |
| 152 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); | 152 | auto CopyL12L0MxScaleA = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleA3510{}); |
| 153 | - AscendC::Te::Copy( | 153 | + CopyL12L0MxScaleA.Call(tensorScaleAL0, tensorBlockScaleAL1, AscendC::Te::MakeCoord(0, kL0Offset)); |
| 154 | - CopyL12L0MxScaleA, tensorScaleAL0, tensorBlockScaleAL1, | ||
| 155 | - AscendC::Te::MakeCoord(0, kL0Offset)); | ||
| 156 | 154 | ||
| 157 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); | 155 | auto layoutBL0 = AscendC::Te::MakeFrameLayout<AscendC::Te::ZNLayoutPtn, AscendC::Std::Int<C0_SIZE>>(curKL0, curN); |
| 158 | auto tensorBL0 = | 156 | auto tensorBL0 = |
| @@ -164,9 +162,7 @@ public: | |||
| 164 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( | 162 | auto tensorScaleBL0 = AscendC::Te::MakeTensor( |
| 165 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); | 163 | AscendC::Te::MakeMemPtr<AscendC::Te::Location::L0B, fp8_e8m0_t>(l0Offset), layoutScaleBL0); |
| 166 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); | 164 | auto CopyL12L0MxScaleB = AscendC::Te::MakeCopy(::Tile::CopyL12L0MxScaleB3510{}); |
| 167 | - AscendC::Te::Copy( | 165 | + CopyL12L0MxScaleB.Call(tensorScaleBL0, tensorBlockScaleBL1, AscendC::Te::MakeCoord(kL0Offset, 0)); |
| 168 | - CopyL12L0MxScaleB, tensorScaleBL0, tensorBlockScaleBL1, | ||
| 169 | - AscendC::Te::MakeCoord(kL0Offset, 0)); | ||
| 170 | 166 | ||
| 171 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 167 | AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| 172 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); | 168 | AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(l0PingPong_ & 0x1); |
| @@ -280,7 +276,7 @@ AscendC::Te::Mmad( | |||
| 280 | } | 276 | } |
| 281 | 277 | ||
| 282 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); | 278 | auto CopyL0C2GM = AscendC::Te::MakeCopy(AscendC::Te::CopyL0C2GM{}); |
| 283 | - AscendC::Te::Copy(CopyL0C2GM, gmC, tensorL0C, AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}); | 279 | + AscendC::Te::Copy(CopyL0C2GM.with(AscendC::Te::FixpipeParams{FINAL_ACCUMULATION}), gmC, tensorL0C); |
| 284 | if (enableL0cPingPong_) { | 280 | if (enableL0cPingPong_) { |
| 285 | l0cPingPong_++; | 281 | l0cPingPong_++; |
| 286 | } | 282 | } |
| @@ -43,5 +43,6 @@ add_dependencies(cann_samples_tensor_api cann_samples_tensor_api_dependencies) | |||
| 43 | target_include_directories(cann_samples_tensor_api INTERFACE | 43 | target_include_directories(cann_samples_tensor_api INTERFACE |
| 44 | "${TENSOR_API_PATH}/include" | 44 | "${TENSOR_API_PATH}/include" |
| 45 | "${TENSOR_API_PATH}/include/tensor_api" | 45 | "${TENSOR_API_PATH}/include/tensor_api" |
| 46 | + "${TENSOR_API_PATH}/include/tensor_api/include" | ||
| 46 | "${ASCEND_DIR}/asc" | 47 | "${ASCEND_DIR}/asc" |
| 47 | ) | 48 | ) |
代码中已将 AscendC::GetTaskRation() 修改为 AscendC::GetTaskRatio()，但 README 示例和说明仍保留旧接口名：
文档与代码不一致会误导用户复制旧接口编译失败。请同步更新 README 中的 GetTaskRation 为 GetTaskRatio。