Huawei Cloud Ascendc Operator Performance Optim

Huawei Cloud AscendC Operator Performance Optimization

Overview

This skill provides guidance for developing and optimizing custom operators using AscendC programming language.

Architecture: Performance Analysis → Bottleneck Identification → Operator Development → Optimization → Validation

Related Skills:

• huawei-cloud-ascend-profiler-db-explorer - Performance data analysis and bottleneck identification
• huawei-cloud-ascend-small-model-migrate - Migration workflow that may require operator optimization

Architecture Components

This skill involves the following cloud services and components:

• AscendC: Programming language for custom operator development
• CANN: Huawei Cloud AI Computing Platform for NPU
• Ascend 910B: Target NPU hardware for operator deployment
• Ascend Profiler: Performance analysis tool for validation

Architecture Diagram:

┌─────────────────────────────────────────────────────────────┐
│            AscendC Operator Optimization Skill             │
├─────────────────────────────────────────────────────────────┤
│  ┌──────────────┐    ┌──────────────┐    ┌──────────────┐ │
│  │  Performance │───▶│  Bottleneck  │───▶│  Operator    │ │
│  │  Analysis    │    │  Identification│   │  Development │ │
│  └──────────────┘    └──────────────┘    └──────────────┘ │
│         │                   │                   │          │
│         ▼                   ▼                   ▼          │
│  ┌──────────────┐    ┌──────────────┐    ┌──────────────┐ │
│  │  Profiling   │    │  Optimization│    │  Validation │ │
│  │  Data        │    │  Techniques  │    │  & Testing  │ │
│  └──────────────┘    └──────────────┘    └──────────────┘ │
└─────────────────────────────────────────────────────────────┘

Use Cases

Typical Problem Scenarios:

• Optimizing performance-critical operators on Ascend NPU
• Developing custom operators for specific workloads
• Improving model inference performance through operator optimization
• Fixing operator bottlenecks identified during profiling
• Implementing missing operators for NPU deployment

Typical User Phrases:

• "Optimize my custom operator for Ascend"
• "Develop AscendC operator for GEMM"
• "Improve inference performance on NPU"
• "Fix bottleneck operator"
• "Implement custom operator using AscendC"
• "AscendCOperator"
• "OptimizationAscendOperatorPerformance"
• "OperatorPerformance"

Scope

Supported:

• Custom operator development in AscendC
• Performance optimization for existing operators
• Operator validation and testing

Not supported:

• Non-AscendC operator development
• Framework-level optimizations

Core Workflow

1. Performance Analysis

• Use profiling tools to identify performance bottlenecks
• Analyze operator execution time and resource utilization

2. Bottleneck Identification

• Identify operators with high execution time
• Determine optimization opportunities

3. Operator Development

• Implement custom operators using AscendC
• Follow AscendC best practices

4. Optimization Techniques

• Memory optimization
• Compute optimization
• Data layout optimization

5. Validation

• Verify functional correctness
• Validate performance improvement

Reference Documents

Prerequisites

• CANN >= 7.0.0 installed
• AscendC >= 1.0.0 installed
• Ascend NPU driver installed and working properly
• Operator code or performance data to be optimized

Core Commands

# Analyze operator performance bottlenecks
msprof --output=/path/to/output ./my_operator

# Optimize operator using AscendC
# Refer to CANN development guide for operator development

Parameter Confirmation

Output Format

Performance analysis results are saved in the specified output directory:

output/
├── summary.json           # Performance summary
├── operator_stats.csv     # Operator execution statistics
├── timeline.json          # Execution timeline data
└── recommendations.md     # Optimization recommendations

Summary JSON Structure:

{
  "total_time_ms": 1234.56,
  "operator_count": 42,
  "top_operators": [
    {"name": "CustomGEMM", "time_ms": 456.78, "percentage": 37.0},
    {"name": "VectorAdd", "time_ms": 123.45, "percentage": 10.0}
  ],
  "optimization_candidates": ["CustomGEMM", "DataTransfer"]
}

Validation Method

Functional Validation

1. Run operator with test inputs 2. Compare outputs with reference implementation 3. Verify numerical accuracy (tolerance: 1e-5 for FP32, 1e-3 for FP16)

Performance Validation

1. Benchmark operator before optimization 2. Apply optimization changes 3. Benchmark operator after optimization 4. Calculate speedup ratio: speedup = time_before / time_after

Acceptance Criteria

• Functional correctness: Output matches reference within tolerance
• Performance improvement: Speedup >= 1.2x (20% improvement)
• No regression: Other operators not affected

Best Practices

Memory Optimization

• Use GM (Global Memory) for large tensors
• Use L1/L0A/L0B for intermediate results in matrix operations
• Align memory access to 32-byte boundaries
• Reuse memory buffers when possible

Compute Optimization

• Vectorize operations using AscendC intrinsics
• Use MMA (Matrix Multiply Accumulate) for matrix operations
• Parallelize independent operations
• Minimize synchronization points

Data Layout Optimization

• Use NZ format for matrix operations
• Use ND format for vector operations
• Avoid unnecessary format conversions
• Consider memory coalescing for data access

Code Structure

• Separate compute logic from memory operations
• Use template metaprogramming for flexibility
• Document optimization assumptions
• Profile before and after each optimization

Notes

Common Pitfalls

• Memory bank conflicts: Ensure data is distributed across memory banks
• Unaligned access: Check 32-byte alignment for all buffers
• Excessive synchronization: Minimize barrier usage between kernels
• Wrong data format: Match format to operation type (NZ for matmul, ND for vector)

Performance Tips

1. Profile first to identify real bottlenecks 2. Focus on hot paths (operators with >10% total time) 3. Consider algorithmic changes before micro-optimizations 4. Test with realistic input sizes 5. Validate correctness after each optimization

Debugging Tips

• Use ASCENDC_DEBUG=1 for verbose logging
• Check CANN log files in /var/log/npu/
• Compare with CPU reference implementation
• Use msprof for detailed performance breakdown

Limitations

• AscendC operators are hardware-specific (910B)
• Not all PyTorch operators have AscendC equivalents
• Custom operators require CANN recompilation for deployment

Acceptance Criteria

Functional Acceptance Criteria

1. Phase 1: Investigation

2. Phase 2: Baseline

3. Phase 3: Optimization

4. Phase 4: Accuracy Verification

5. Phase 5: Performance Verification

Correct/Error Pattern Comparison

Directory Backup

Correct: Backup before modification

# Backup with timestamp
cp -r operator_dir operator_dir_backup_$(date +%Y%m%d%H%M%S)

# Verify backup
ls -la operator_dir_backup_*/

Error: Modify without backup

# Direct modification is risky
vi op_kernel/add.cpp  # Lost original if something goes wrong

Anti-Pattern Violations

Correct: Cast FP16/BF16 to FP32 for complex math

half a = ...;
float a_f = (float)a;
// Use a_f for sqrt/exp/log etc.
float result = expf(a_f);
half result_h = (half)result;

Error: Direct FP16/BF16 math

half a = ...;
// Wrong: exp expects float
half result = exp(a);  // Undefined behavior

Compilation

Correct: Use provided build scripts

bash build.sh
# Or
mkdir build && cd build
cmake .. && make

Error: Manual compilation without cmake

g++ -o op op.cpp  # Missing include paths, defines

Non-Functional Acceptance Criteria

Test Cases Summary

Positive Test Cases

1. TC-001: Tiling optimization 2. TC-002: Data copy optimization 3. TC-003: API usage optimization 4. TC-004: Memory optimization 5. TC-005: Pipeline optimization 6. TC-006: End-to-end optimization workflow

Negative Test Cases

1. TC-N01: Modify without backup 2. TC-N02: Violate anti-pattern rules 3. TC-N03: Compilation failure 4. TC-N04: Accuracy degradation 5. TC-N05: Performance not improved

Phase 3: API UsageOptimization — DetailedReference

3.1 TPipe in kernel Create Outside Class

TPipe As Class Membertime, InitializationableSetGlobal TPipe fingerneedle, CompileadapterrecognizeastypeMemoryemptybetweenhavebe externalpartdirtydyeofRisk, causethisreleaseabandonpairtypeinside Scalar changeamountofoftenamountfoldstackandoftenamounttransferbroadcastOptimization.

reverseexample — TPipe asClass member:

template <typename ComputeT> class KernelExample {
public:
    __aicore__ inline KernelExample() {}
    __aicore__ inline void Init(...) {
        pipe.InitBuffer(xxxBuf, BUFFER_NUM, xxxSize);
    }
private:
    TPipe pipe;       // ← intypeInternal, blockstop Scalar Optimization
};

extern "C" __global__ __aicore__ void example_kernel(...) {
    KernelExample<float> op;
    op.Init(...);
}

positiveexample — TPipe inCreate Outside Class, in order tofingerneedletransferinput:

template <typename ComputeT> class KernelExample {
public:
    __aicore__ inline KernelExample() {}
    __aicore__ inline void Init(..., TPipe* pipeIn) {
        pipe = pipeIn;
        pipe->InitBuffer(xxxBuf, BUFFER_NUM, xxxSize);
    }
private:
    TPipe* pipe;      // ← onlykeepfingerneedle, typeMemoryemptybetweeninterfereclean
};

extern "C" __global__ __aicore__ void example_kernel(...) {
    TPipe pipe;                     // ← inCreate Outside Class
    KernelExample<float> op;
    op.Init(..., &pipe);
}

actualtest: averageaverage scalar_time from 281 us downgradearrive 236 us (−17%) , scalar_time occupycomparefrom 21% downgradearrive 17%. anywhatScenariosallSuggestUsagethisOptimization, scalar bound Scenariosreceiveadvantageousespeciallyascleardisplay.

3.2 pureData CopyOperatorUsage TQueBind

pureData CopyOperatornotinvolveand Vector Calculation, standardstandard VECIN→VECOUT modelformulaableleadinputonetimeredundantremainingof LocalTensor→LocalTensor DataCopy.

reverseexample — redundantremainingof Vector copyshell:

TQue<QuePosition::VECIN, BUFFER_NUM> QueI;
TQue<QuePosition::VECOUT, BUFFER_NUM> QueO;

auto iLocal = QueI.AllocTensor<ComputeT>();
DataCopy(iLocal, inGm[i * 32], size);
QueI.EnQue(iLocal);
auto iLocal2 = QueI.DeQue<ComputeT>();
for (int j = 0; j < jLen; ++j) {
    auto oLocal = QueO.AllocTensor<ComputeT>();
    DataCopy(oLocal, iLocal2, size);   // LocalTensor → LocalTensor, wavecost Vector
    QueO.EnQue(oLocal);
    auto oLocal2 = QueO.DeQue<ComputeT>();
    DataCopyPad(outGm[j], oLocal2, ...);
    QueO.FreeTensor(oLocal2);
}
QueI.FreeTensor(iLocal2);

positiveexample — TQueBind Eliminateredundantremainingcopyshell:

TQueBind<QuePosition::VECIN, QuePosition::VECOUT, BUFFER_NUM> queBind;

auto bindLocal = queBind.AllocTensor<ComputeT>();
DataCopy(bindLocal, inGm[i * 32], size);
queBind.EnQue(bindLocal);
auto bindLocal2 = queBind.DeQue<ComputeT>();
for (int j = 0; j < len; ++j) {
    DataCopyPad(outGm[j], bindLocal2, ...);
}
queBind.FreeTensor(bindLocal2);

validresult: aiv_vec_time downgradearriveapproximately 0.

3.3 Counter modelformula (SetMaskCount)

Normal modelformularequireshandmoveCalculationmainblock/tailblockof mask anditeraterepresenttimenumber, involveandlargeamount Scalar openconsume. Counter modelformulastraightconnecttransferinputtotalunitelementnumber, HardwareAutomaticpushjudgeiteraterepresenttimenumber.

reverseexample — Normal modelformula (half typetype, 15000 countunitelement) :

uint32_t ELE_SIZE = 15000;
AscendC::BinaryRepeatParams binaryParams;
uint32_t numPerRepeat    = 256 / sizeof(DTYPE_X);   // half → 128
uint32_t mainRepeatTimes = ELE_SIZE / numPerRepeat;  // 117
uint32_t tailEleNum      = ELE_SIZE % numPerRepeat;  // 24

AscendC::SetMaskNorm();
AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::NORMAL>(numPerRepeat);
AscendC::Add<DTYPE_X, false>(zLocal, xLocal, yLocal,
    AscendC::MASK_PLACEHOLDER, mainRepeatTimes, binaryParams);
if (tailEleNum > 0) {
    AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::NORMAL>(tailEleNum);
    AscendC::Add<DTYPE_X, false>(
        zLocal[mainRepeatTimes * numPerRepeat],
        xLocal[mainRepeatTimes * numPerRepeat],
        yLocal[mainRepeatTimes * numPerRepeat],
        AscendC::MASK_PLACEHOLDER, 1, binaryParams);
}
AscendC::ResetMask();

positiveexample — Counter modelformula, onetimeadjustuse:

uint32_t ELE_SIZE = 15000;
AscendC::BinaryRepeatParams binaryParams;
AscendC::SetMaskCount();
AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::COUNTER>(ELE_SIZE);
AscendC::Add<DTYPE_X, false>(zLocal, xLocal, yLocal,
    AscendC::MASK_PLACEHOLDER, 1, binaryParams);
AscendC::ResetMask();

whenmultiplecount Vector fingercommandhandlemanagemutualsameunitelementnumberamounttime, Counter modelformulaoptimizetrendupdatecleardisplay——noneedreversecomplex CalculationDifferentmainblock/tailblock mask.

3.4 Matmul AtomicAdd

Matmul Result C(m,n) requiresand GM aboverulematrix D(m,n) mutualaddtime, canintransferoutputPathabovemergematch.

reverseexample — handmoveData Copyafterin UB do Add:

mm.IterateAll(local_c);

DataCopy(local_d, gm_d, d_size);
event_t eventId = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
SetFlag<HardEvent::MTE2_V>(eventId);
WaitFlag<HardEvent::MTE2_V>(eventId);
Add(local_d, local_d, local_c, d_size);
DataCopy(gm_d, local_d, d_size);

positiveexample — AtomicAdd mergematchtotransferoutput:

mm.IterateAll(gm_d, 1);          // enAtomic = 1
// orin Iterate loopringmiddle: 
// mm.GetTensorC(gm_d, 1);       // enAtomic = 1

M=64, N=256, K=256 actualtest: averageaverage cycle from 154181 downgradearrive 135054 (−12.4%) .

3.5 returnapproximatelyfingercommandgroupmatch

willconnectcontinue buffer allpartaccumulateaddasonecountstandardamountofScenarios:

pushrecommendmodelformula (float, shape=256) :

static constexpr uint32_t BLK_LEN = 32;
TBuf<QuePosition::VECCALC> calcBuf;
pipe.InitBuffer(calcBuf, totalLength * sizeof(float));
AscendC::LocalTensor<float> tempTensor1 = calcBuf.Get<float>();
constexpr uint32_t c0Count = BLK_LEN / sizeof(float);
const uint32_t blockNum0 = (totalLength + c0Count - 1) / c0Count;

AscendC::SetMaskCount();
AscendC::SetVectorMask<float>(0, totalLength);
AscendC::BlockReduceSum<float, false>(tempTensor1, xLocal,
    AscendC::MASK_PLACEHOLDER, 1,
    DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();

AscendC::SetVectorMask<float>(0, blockNum0);
AscendC::WholeReduceSum<float, false>(zLocal, tempTensor1,
    AscendC::MASK_PLACEHOLDER, 1,
    DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetMaskNorm();

AscendC BaseDataStructureInterfaceSummary

one, LocalTensor

Purpose: StoreAI CoreInternalLocal MemoryData, Logical LocationPackageincludeVECIN, VECOUT, VECCALC, A1, A2, B1, B2, CO1, CO2.

structurecreateandInitialization

// PipeFramework (notstraightconnectadjustuse) 
AscendC::LocalTensor<T>() {}

// quietstateTensorcompileprocess
AscendC::LocalTensor<T>(TPosition pos, uint32_t addr, uint32_t tileSize)
AscendC::LocalTensor<T>(uint32_t addr)  // onlySupportTensorTraittypetype

CoreInterface

Examples

// distributematchandUsage
AscendC::LocalTensor<half> srcLocal = inQueue.AllocTensor<half>();
AscendC::DataCopy(srcLocal, srcGlobal, 512);
inQueue.EnQue(srcLocal);

// unitelementvisitask
srcLocal.SetValue(0, 1.0f);
auto val = srcLocal.GetValue(0);

// partialmoveoperatework
AscendC::LocalTensor<half> offsetTensor = srcLocal[16];

// typetypeconvertexchange
AscendC::LocalTensor<int16_t> castTensor = srcLocal.ReinterpretCast<int16_t>();

---

two, GlobalTensor

Purpose: StoreGlobal MemoryGlobalData.

CoreInterface

Examples

AscendC::GlobalTensor<half> srcGlobal;
srcGlobal.SetGlobalBuffer((__gm__ half*)srcGm, dataSize);

// Read
auto val = srcGlobal.GetValue(0);

// partialmovevisitask
AscendC::GlobalTensor<half> offsetGlobal = srcGlobal[128];

// DataCopy
AscendC::DataCopy(srcLocal, srcGlobal, dataSize);

---

three, Layout

Purpose: DescriptionmultipledimensionexpandamountMemoryarrangelocal, PackagecontainShapeandStride.

reasontype

template <typename ShapeType, typename StrideType>
struct Layout {
    __aicore__ inline constexpr Layout(const ShapeType& shape = {}, const StrideType& stride = {});
    __aicore__ inline constexpr decltype(auto) GetShape();
    __aicore__ inline constexpr decltype(auto) GetStride();
    template <typename CoordType>
    __aicore__ inline constexpr auto operator()(const CoordType& coord) const;
};

structurecreateMethod

#include "kernel_operator_layout.h"

// Shapestructurecreate
auto shape = AscendC::MakeShape(4, 2);  // 4travel2column

// Stridestructurecreate
auto stride = AscendC::MakeStride(4, 1);  // travelsteplength4, columnsteplength1

// Layoutstructurecreate
auto layout = AscendC::MakeLayout(shape, stride);

// throughexceedsitstandardCalculationMemorysearchlead
auto coord = AscendC::MakeCoord(1, 0);  // chapter1travelchapter0column
auto idx = layout(coord);  // Calculationobtaintoregionaddresssearchlead

Examples: 4travel2columnrulematrix

Shape: (4, 2), Stride: (4, 1)

---

four, Coordinate

Purpose: tableshowexpandamountmultipledimensionsitstandard, matchmatchLayoutUsageCalculationMemorysearchlead.

reasontype

template <typename... Coords>
using Coord = Std::tuple<Coords...>;

Interface

// structurecreatesitstandard
auto coord = AscendC::MakeCoord(row, col);

// sitstandardconvertMemorysearchlead
template <typename CoordType, typename ShapeType, typename StrideType>
__aicore__ inline constexpr auto Crd2Idx(const CoordType& coord,
                                         const Layout<ShapeType, StrideType>& layout);

Examples

auto shape = AscendC::MakeShape(4, 2);
auto stride = AscendC::MakeStride(4, 1);
auto layout = AscendC::MakeLayout(shape, stride);

auto coord = AscendC::MakeCoord(1, 0);  // row=1, col=0
auto idx = AscendC::Crd2Idx(coord, layout);  // idx = 4

---

five, TensorTrait

Purpose: DescriptionTensorofcompleteadjustinformationinformation (Datatypetype, Logical Location, Layout) , Used forCompileperiodOptimization.

reasontype

template <typename T, TPosition pos = TPosition::GM,
          typename LayoutType = Layout<Shape<>, Stride<>>>
struct TensorTrait {
    using LiteType = T;
    using LiteLayoutType = LayoutType;
    static constexpr const TPosition tPos = pos;

    __aicore__ inline LayoutType& GetLayout();
    __aicore__ inline void SetLayout(const LayoutType& t);
};

structurecreateMethod

#include "kernel_operator_tensor_trait.h"

auto shape = AscendC::MakeShape(16, 16, 16);
auto stride = AscendC::MakeStride(0, 0, 0);
auto layout = AscendC::MakeLayout(shape, stride);

// structurecreateTensorTrait
auto tensorTrait = AscendC::MakeTensorTrait<float, AscendC::TPosition::VECIN>(layout);

// UsageTensorTraittypetypestructurecreateLocalTensor
AscendC::LocalTensor<decltype(tensorTrait)> tensor(addr);

SupportofDatatypetype

int4b_t, uint8_t, int8_t, int16_t, uint16_t, bfloat16_t, int32_t, uint32_t, int64_t, uint64_t, float, half

Constraints

• sameoneInterfacenotSupportsametimeoutputinputTensorTraittypetypeandnonTensorTraittypetypeofTensor
• TensorTraittypetypeofTensorDoes Not IncludeShapeInfoinformationinformation
• DataCopycutsliceInterfacenotSupportTensorTraittypetype

---

TPosition Logical Location

AscendC DataData Copy API Reference

⚠️ Production Rules

GM ↔ UB Data Copymust use DataCopyPad, Do Not Use DataCopy.

DataCopyPad (pushrecommend)

GM → UB

AscendC::DataCopyExtParams copyParams{blockCount, blockLen, srcStride, dstStride, 0};
AscendC::DataCopyPadExtParams<T> padParams{isPad, leftPad, rightPad, padValue};
AscendC::DataCopyPad(dstLocal, srcGlobal, copyParams, padParams);

UB → GM

AscendC::DataCopyExtParams copyParams{blockCount, blockLen, srcStride, dstStride, 0};
AscendC::DataCopyPad(dstGlobal, srcLocal, copyParams);

DataCopyExtParams parameternumber

⚠️ Stride singlebitnotsame: GM sideusecharactersection, UB sideuse 32B DataBlock. thisismostoftenseeofData Copy bug Source.

DataCopyPadExtParams parameternumber (only GM→UB)

rLength vs rLengthAlign usemethod

relatedkey: CopyOut of srcStride isblockbetweenbetweenseparate (padding partdistribute) , notiscompleteadjusttravellengthdegree.

CopyIn/CopyOut oneconsistentproperty

CopyIn use DataCopyPad time, CopyOut alsomustmustuse DataCopyPad. mixuseableguideconsistenttravelerrorbit.

BaseDataData Copy (DataCopy)

onlyUsed for UB ↔ UB Internalcopyshell.

// UB -> UB
AscendC::DataCopy(dstLocal, srcLocal, count);

parameternumber:

• count: unitelementcountnumber, count * sizeof(T) need32charactersectionpairalign

nonconnectcontinueData Copy (DataCopyParams)

AscendC::DataCopyParams params;
params.blockCount = 1;   // connectcontinueDatablockcountnumber [1, 4095]
params.blockLen = 8;     // eachblocklengthdegree, singlebitDataBlock(32B) [1, 65535]
params.srcGap = 0;       // sourcemutualneighborblockbetweenseparate, singlebitDataBlock(32B)
params.dstGap = 0;       // itemofmutualneighborblockbetweenseparate, singlebitDataBlock(32B)

AscendC::DataCopy(dstLocal, srcGlobal, params);

showmeaningfigure:

blockCount=2, blockLen=8, srcGap=0, dstGap=1
source: [====8block====][====8block====]
itemof: [====8block====][gap][====8block====]

cutsliceDataData Copy (SliceInfo)

AscendC::SliceInfo srcSliceInfo[] = {{16, 70, 7, 3, 87}, {0, 2, 1, 1, 3}};
AscendC::SliceInfo dstSliceInfo[] = {{0, 47, 0, 3, 48}, {0, 1, 0, 1, 2}};
uint32_t dimValue = 2;

AscendC::DataCopy(dstLocal, srcGlobal, dstSliceInfo, srcSliceInfo, dimValue);

SliceInfo Structure:

nonpairalignData Copy (DataCopyPad)

// GM -> UB, Supportnon32charactersectionpairalign
AscendC::DataCopyExtParams copyParams{1, 20 * sizeof(half), 0, 0, 0};
AscendC::DataCopyPadExtParams<half> padParams{true, 0, 2, 0};  // isPad, leftPad, rightPad, padValue
AscendC::DataCopyPad(dstLocal, srcGlobal, copyParams, padParams);

// UB -> GM
AscendC::DataCopyPad(dstGlobal, srcLocal, copyParams);

DataCopyExtParams:

DataCopyPadExtParams:

UBInternalcopyshell (Copy)

// VECIN/VECCALC/VECOUT ofbetweenofcopyshell
AscendC::Copy(dstLocal, srcLocal, mask, repeatTime, {dstStride, srcStride, dstRepStride, srcRepStride});

CopyRepeatParams:

// Examples: connectcontinuecopyshell512countint16_t
uint64_t mask = 128;
AscendC::Copy(dstLocal, srcLocal, mask, 4, {1, 1, 8, 8});

increasestrongDataData Copy (DataCopyEnhancedParams)

AscendC::DataCopyParams intriParams;
AscendC::DataCopyEnhancedParams enhancedParams;
enhancedParams.blockMode = BlockMode::BLOCK_MODE_MATRIX;  // or BLOCK_MODE_VECTOR
AscendC::DataCopy(dstLocal, srcLocal, intriParams, enhancedParams);

blockMode modelformula:

Quantizationmodelformula (deqScale):

Datathroughpathspeedsearch

regionaddresspairalignRequirements

CommonCodemodelformula

CopyIn (multipletravelbatchamounttransferinput)

__aicore__ inline void CopyIn(uint32_t startRow, uint32_t rows) {
    LocalTensor<half> srcLocal = inQueue.AllocTensor<half>();
    // blockCount=travelnumber, blockLen=eachtravelhavevalidcharactersection, srcStride=GMtravelbetweendistance(charactersection), dstStride=UBtravelbetweendistance(32Bblock)
    AscendC::DataCopyExtParams copyParams{
        static_cast<uint16_t>(rows),                          // blockCount
        static_cast<uint32_t>(cols * sizeof(half)),           // blockLen (havevalidData)
        static_cast<uint32_t>(totalCols * sizeof(half)),      // srcStride (GM, charactersection)
        static_cast<uint16_t>(alignedCols * sizeof(half) / 32) // dstStride (UB, 32Bblock)
    };
    AscendC::DataCopyPadExtParams<half> padParams{true, 0,
        static_cast<uint8_t>(alignedCols - cols), 0};
    AscendC::DataCopyPad(srcLocal, srcGlobal[startRow * totalCols], copyParams, padParams);
    inQueue.EnQue(srcLocal);
}

CopyOut (multipletravelbatchamounttransferoutput)

__aicore__ inline void CopyOut(uint32_t startRow, uint32_t rows) {
    LocalTensor<half> dstLocal = outQueue.DeQue<half>();
    AscendC::DataCopyExtParams copyParams{
        static_cast<uint16_t>(rows),
        static_cast<uint32_t>(cols * sizeof(half)),            // onlytransferoutputhavevalidData
        static_cast<uint16_t>((alignedCols - cols) * sizeof(half) / 32), // srcStride: padding betweenseparate
        static_cast<uint32_t>(totalCols * sizeof(half))        // dstStride (GM, charactersection)
    };
    AscendC::DataCopyPad(dstGlobal[startRow * totalCols], dstLocal, copyParams);
    outQueue.FreeTensor(dstLocal);
}

Elementwise connectcontinueData Copy

__aicore__ inline void CopyIn() {
    LocalTensor<half> srcLocal = inQueue.AllocTensor<half>();
    AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(tileLength * sizeof(half)), 0, 0, 0};
    AscendC::DataCopyPad(srcLocal, srcGlobal[offset], copyParams);
    inQueue.EnQue(srcLocal);
}

regionaddresspairalignRequirements

32 charactersectionpairalignCalculation

// according to 32B pairalignofunitelementnumber
uint32_t alignedCols = ((cols * sizeof(T) + 31) / 32) * (32 / sizeof(T));

// etcvalueFormula
uint32_t elemsPerBlock = 32 / sizeof(T);  // half:16, float:8
uint32_t alignedCols = ((cols + elemsPerBlock - 1) / elemsPerBlock) * elemsPerBlock;

AscendC OperatorCode GenerationReference DocumentsLoadGuide

thisDocumentsfingerguide agent inDifferentOperatorDevelopmentScenariosunder, Load on DemandCorrespondingof reference Documents, avoidavoidonetimepropertyLoadallpartDocumentscreatebecomeaboveunderdocumentwavecost.

ReferencesDocumentsChecklist

ScenariostransformLoadstrategystrategy

Scenarios 1: Elementwise Operator (ReLU, GELU, Add, Mul etc)

specialfeature: Iterativeunitelementoperatework, outputinputOutput shape mutualsame, onedimensionTiling

mustmustLoad:

• basic-data-structures-api.md — GlobalTensor/LocalTensor usemethod
• resource-management-api.md — TPipe/TQue/TBuf Initialization, Double Buffer
• data-copy-api.md — DataCopyPad connectcontinueData Copy
• vector-compute-api.md — computeart/oneunit/standardamountruncompute, Cast upgradeprecisiondegreemodelformula
• kernel-constraints.md — Prohibited std::, repeatTime Limitations

notrequiresLoad:

• sync-control-api.md — unitelementlevelOperatornocorebetweenaccordingdepend, EnQue/DeQue alreadyinResourceManagementmiddleDescription

Scenarios 2: returnapproximately/returnonetransformtypeOperator (LayerNorm, Softmax, BatchNorm etc)

specialfeature: Packagecontain ReduceSum/ReduceMax, according totravel/dimensiondegreeTiling, canabilityrequires FP32 middlebetweenprecisiondegree

mustmustLoad:

• basic-data-structures-api.md — GlobalTensor/LocalTensor
• resource-management-api.md — TPipe/TQue/TBuf, UB contentamountCalculation, blockCount Limitations
• data-copy-api.md — DataCopyPad multipletravelData Copy, rLength/rLengthAlign usemethod, Stride singlebit
• vector-compute-api.md — weightpoint: returnapproximately API (Level2/Pattern) , tmpBuffer Calculation, standardamountOptimization (Adds/Muls) , Cast mixmatchprecisiondegree, multipletravelwidebroadcast
• kernel-constraints.md — repeatTime≤255 distributebatch, Compare pairalign

notrequiresLoad:

• sync-control-api.md — travellevelaloneestablishreturnapproximatelynocorebetweenaccordingdepend

Scenarios 3: pooltransformtypeOperator (AvgPool, MaxPool etc)

specialfeature: slipperymovewindowmouthoperatework, multipledimensiondegreeiteratehistory, coreinsidehavecomplexmixedloopringStructure

mustmustLoad:

• basic-data-structures-api.md — GlobalTensor/LocalTensor
• resource-management-api.md — TPipe/TQue/TBuf, accumBuf etcmultipleapproachtimeslowconflictregion
• data-copy-api.md — multipletime DataCopyPad transferinputnotsametravel/bitplace
• vector-compute-api.md — accumulateadd, typetypeconvertexchange, Duplicate Initialization
• kernel-constraints.md — throughuse Kernel Limitations

notrequiresLoad:

• sync-control-api.md — each slice aloneestablishhandlemanage

Scenarios 4: requirescorebetweenSynchronizationofOperator (AllReduce, Globalreturnapproximatelyetc)

specialfeature: Multi-coreofbetweenkeepinDataaccordingdepend, requiresfirstlocalpartCalculationagainGlobalMerge

allpartLoad:

• basic-data-structures-api.md
• resource-management-api.md — requires Workspace Management (GM + UB workspace)
• data-copy-api.md
• vector-compute-api.md
• sync-control-api.md — weightpoint: SyncAll/IBSet/IBWait corebetweenSynchronization, workspace emptybetweenRequirements
• kernel-constraints.md

Scenarios 5: onlymodifymodifyalreadyhaveOperator (bug modifycomplex, smallrangescopemodifymove)

Load on Demand: onlyLoadandmodifymodifyRelatedofDocuments. exampleif:

• modifymodifyCalculationlogiclogic → vector-compute-api.md + kernel-constraints.md
• modifymodifyDataData Copy → data-copy-api.md
• modifymodifyResourcedistributematch → resource-management-api.md
• Runtime crash / Dataerrorerror → optimizefirst kernel-constraints.md diagnosejudgeChecklist

Scenarios 6: PerformanceOptimization

mustmustLoad:

• resource-management-api.md — Double Buffer reasonmanage, UB utilizeuserate
• vector-compute-api.md — standardamountOptimization (Adds/Muls representsubstitute Duplicate+runcompute) , multipletravelwidebroadcast, Pattern returnapproximately
• kernel-constraints.md — repeatTime distributebatchOptimization

throughuseRules

1. mostsmallLoadreasonrule: optimizefirstLoadandwhenpreviousOperatortypetypestraightconnectRelatedofDocuments, avoidavoidnorelatedDocumentsdisappearconsumeaboveunderdocument 2. LimitationsDocumentsoftenLoad: kernel-constraints.md Packagecontainhighfrequencystamppitfallpoint, newOperatorDevelopmenttimeSuggestbeginfinalLoad 3. according tocompilecodePhaseLoad: compilecompose Init timesideweight resource-management-api.md, compilecompose Compute timesideweight vector-compute-api.md 4. meettoCompile/Runerrorerrortime: optimizefirstsearchsee kernel-constraints.md diagnosejudgeChecklist 5. DataData Copyaskproblem: optimizefirstCheck data-copy-api.md middleof Stride singlebitand rLength/rLengthAlign regiondistribute

AscendC Kernel LimitationsandPitfallsGuide

ProhibitedUsageof C++ Features (Kernel side)

Standard LibraryMath Functions

Kernel CodemiddleProhibitedUsage std:: commandnameemptybetweenunderofMath Functions, CompiletimecanabilitythroughexceedbutRuntimeproduceliveerrorerrorResult:

movestateMemorydistributematch

Kernel middleProhibitedUsageanywhatmovestateMemorydistributematch:

Host/Kernel headfileseparateleave

repeatTime overflowoutput

allhaveUsagehighdimensionTilingmodelformulaof API (Add, Sub, Mul, Div, Cast, Duplicate etc) , other repeatTime parameternumbertypetypeas `uint8_t`, mostlargevalue 255.

quietsilentcutjudge: transferinput 256 ablebecutjudgeas 0, guideconsistentnotExecuteanywhatCalculationandnoreporterror.

Host Sideguardshield

// Tiling PhaseLimitationsmostlargetravelnumber
tileRows = std::min(tileRows, static_cast<uint32_t>(255));

Kernel Sidedistributebatch

int64_t remaining = rowCount;
int64_t offset = 0;
while (remaining > 0) {
    uint8_t batch = static_cast<uint8_t>(std::min(remaining, (int64_t)255));
    AscendC::Sub(dst[offset], src0[offset], src1, mask, batch, params);
    offset += batch * alignedCols;
    remaining -= batch;
}

receiveshadowloudof API

allhaveconnectreceive repeatTime parameternumberofhighdimensionTilingweightload: Add, Sub, Mul, Div, Adds, Muls, Cast, Duplicate, Compare, Select, Exp, Ln, Abs, Sqrt, Reciprocal etc.

Compare API 256 charactersectionpairalign

Compare RequirementsparticipationcomparecompareofDataregiondomainas 256 charactersectionadjustnumbertimes. notfootpartdistributeneed padding:

• ArgMax → padding fill -inf or -FLT_MAX
• ArgMin → padding fill +inf or FLT_MAX

uint32_t align256Elems = 256 / sizeof(T);
uint32_t alignedCount = ((count + align256Elems - 1) / align256Elems) * align256Elems;
if (alignedCount > count) {
    AscendC::Duplicate(src[count], paddingValue, alignedCount - count);
}

oftenamountandCompileperiodOptimization

• optimizefirstUsage constexpr definemeaningCompileperiodoftenamount
• avoidavoidRuntimeCalculationcanin order toinCompileperiodDetermineofvalue
• 32 charactersectionpairalignCalculation: ((x + 31) / 32) * 32

API blacknamesingle

onlyallowallowadjusttrytimeUsage:

AscendC::printf("debug: xGm[0]=%f\n", xGm.GetValue(0));  // onlyadjusttry

diagnosejudgeCheckChecklist

meetto Kernel CompileorRunerrorerrortimeaccording tothissequenceorderarrangesearch:

1. iswhetherUsagecompleted std:: functionnumber? → substituteexchangeas AscendC API 2. DataData Copyiswhetherusecompleted DataCopyPad? → GM↔UB mustmustuse DataCopyPad 3. repeatTime iswhetherexceedexceed 255? → distributebatchhandlemanage 4. Compare Dataiswhether 256B pairalign? → padding 5. ReduceMax/Sum of dst and tmpBuffer iswhethernotsame? → distributeopendistributematch 6. InitBuffer totalnumberiswhetherexceedexceed 64? → Merge buffer 7. EnQue/DeQue iswhethermatchpair? → Data CopyaftermustmustSynchronization

AscendC ResourceManagementInterfaceSummary

one, TPipe

Purpose: statisticsoneManagementDeviceSideMemory andSynchronizationmatterfileResource, onecountKernelmustmustandonlyabilityhaveonecountTPipepairimage.

Corefunctionability

1. MemoryResourceManagement: throughexceedInitBufferasTQueandTBufdistributematchMemory 2. SynchronizationmatterfileManagement: throughexceedAllocEventID/ReleaseEventIDManagementmatterfileID

relatedkeyInterface

Examples

AscendC::TPipe pipe;

// asTQuedistributematchMemory
AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue;
pipe.InitBuffer(inQueue, 2, 1024);  // 2block, eachblock1024charactersection

// asTBufdistributematchMemory
AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuf;
pipe.InitBuffer(tmpBuf, 512);  // 512charactersection

// EventIDManagement
AscendC::TEventID eventId = GetTPipePtr()->AllocEventID<AscendC::HardEvent::V_S>();
// ... UsageEventID ...
GetTPipePtr()->ReleaseEventID<AscendC::HardEvent::V_S>(eventId);

---

two, TQue

Purpose: Managementteamcolumnoperatework, ImplementationPipelinelineParallel.

Templateparameternumber

template <TPosition pos, int32_t depth, auto mask = 0>
class TQue {...};

CoreInterface

BuffernumberamountLimitations

standardstandardPipelinelinemodelformula

// CopyIn -> Compute -> CopyOut
AscendC::LocalTensor<half> srcLocal = inQueue.AllocTensor<half>();
AscendC::DataCopy(srcLocal, srcGlobal, dataSize);
inQueue.EnQue(srcLocal);

srcLocal = inQueue.DeQue<half>();
// ... Calculationoperatework ...
inQueue.FreeTensor(srcLocal);

Double Buffer modelformula

Double Buffer ofitemofislet DMA Data Copy (MTE2/MTE3) and Vector CalculationParallelExecute, andnonsimplesingleof"twoblockMemorydoCalculation".

// MTE throughpathteamcolumn: depth=1, num=2 (double buffer)
AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue;
AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue;
pipe.InitBuffer(inQueue, 2, len);   // 2block: oneblockinData Copy, oneblockinCalculation
pipe.InitBuffer(outQueue, 2, len);

// pureCalculationapproachtimeslowconflictregion: noneed double buffer
AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuf;
pipe.InitBuffer(tmpBuf, len);       // 1blockimmediatecan

whattimeuse TQue vs TBuf:

---

three, TBuf

Purpose: ManagementapproachtimechangeamountMemory, notSupportinputteamoutputteamoperatework.

specialpoint

• onlyabilityparticipationCalculation, nomethodExecuteteamcolumnoperatework
• TPipeonlyasTBufdistributematchoneblockMemory
• obtaingetofTensornoneedexplainrelease
• multiplecountapproachtimechangeamountrequiresdefinemeaningmultiplecountTBuf

Interface

Examples

AscendC::TPipe pipe;
AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuf;
pipe.InitBuffer(tmpBuf, 512);

// obtaingetTensorUsage
AscendC::LocalTensor<float> tmp = tmpBuf.Get<float>();
// UsagetmpperformCalculation, noneedexplainrelease

---

four, TQueBind

Purpose: binddefineVECINandVECOUTImplementationMemorycomplexuse.

Templateparameternumber

template <TPosition srcPos, TPosition dstPos, int32_t depth>
class TQueBind {...};

UsageScenarios

Used forkeepinVectorCalculationtimeImplementationVECINandVECOUTMemorycomplexuse.

Examples

AscendC::TQueBind<AscendC::TPosition::VECIN, AscendC::TPosition::VECOUT, 1> que;
pipe.InitBuffer(que, 2, 1024);

AscendC::LocalTensor<half> tensor = que.AllocTensor<half>();
que.EnQue<AscendC::TPosition::GM, AscendC::TPosition::VECIN, half>(tensor);
tensor = que.DeQue<AscendC::TPosition::GM, AscendC::TPosition::VECIN, half>();

---

five, WorkspaceManagement

GetUserWorkspace

obtaingetuseuserUsageofworkspacefingerneedle:

__aicore__ inline GM_ADDR GetUserWorkspace(GM_ADDR workspace);

GM_ADDR usrWorkspace = AscendC::GetUserWorkspace(workspace);

GetSysWorkSpacePtr

obtaingetsystemstatisticsworkspacefingerneedle (Used forMatmuletchighlevelAPI) :

__aicore__ inline __gm__ uint8_t* GetSysWorkSpacePtr();

REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm, &tiling);

SetSysWorkSpace

Setsystemstatisticsworkspace (KernelstraightadjustScenarios) :

__aicore__ inline void SetSysWorkspace(GM_ADDR workspace);

AscendC::SetSysWorkspace(workspace);
if (GetSysWorkSpacePtr() == nullptr) {
    return;
}

HostsideConfiguration

// Tilingfunctionnumbermiddle
size_t usrSize = 256;
auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
size_t *currentWorkspace = context->GetWorkspaceSizes(1);
currentWorkspace[0] = usrSize + sysWorkspaceSize;

---

six, MemoryManagementConstraints

InitBuffer Constraints

• applypleaseofMemoryablein TPipe analyzestructuretimeAutomaticexplainrelease
• onecount kernel middleallhave Buffer numberamountofandnotabilityexceedexceed 64
• CustomregionaddressmethodsandnotfingerdefineregionaddressmethodsnotSuggestmixuse
• len notfullfoot 32 charactersectionpairaligntimeableAutomaticsupplementalign

UB contentamountandtravelnumberCalculation

Host Sidethroughexceedaverageplatform API obtainget UB largesmallafter, Calculationeachtimehandlemanageoftravelnumber:

// Host side
uint64_t ubSize;
ascendc_platform->GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);

// eachtraveloccupyuse UB charactersection = alignedCols * sizeof(T)
// considerconsidermultiplecount buffer (if inQueue×2 + outQueue×2 + tmpBuf×1 + calcBuf×1 = 6portion) 
uint32_t tileRows = ubSize / (alignedCols * sizeof(T) * bufferCount);

DataCopyPad blockCount Limitations

DataCopyExtParams.blockCount mostlargevalueas 4095. when tileRows > 4095 timerequiresdistributebatch:

tileRows = std::min(tileRows, (uint32_t)4095);

AllocTensorConstraints

sameoneTPositionaboveconnectcontinueAllocofTensornumberamountLimitations:

solveresolveBuffernotfoot

// Method1: Mergemultiplecountbuffertooneblock, throughexceedpartialmoveUsage
pipe.InitBuffer(que0, 1, len * 3);
AscendC::LocalTensor<T> local1 = que0.AllocTensor<T>();
AscendC::LocalTensor<T> local2 = local1[len];
AscendC::LocalTensor<T> local3 = local1[len * 2];

// Method2: explainreleasenotuseofTQue
que0.FreeAllEvent();

---

seven, classictypeUsagemodelformula

VectorOperatorstandardstandardmodelformula

class KernelOp {
public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t size) {
        xGlobal.SetGlobalBuffer((__gm__ half*)x, size);
        yGlobal.SetGlobalBuffer((__gm__ half*)y, size);
        pipe.InitBuffer(inQueue, 1, size * sizeof(half));
        pipe.InitBuffer(outQueue, 1, size * sizeof(half));
    }

    __aicore__ inline void Process() {
        CopyIn();
        Compute();
        CopyOut();
    }

private:
    __aicore__ inline void CopyIn() {
        AscendC::LocalTensor<half> xLocal = inQueue.AllocTensor<half>();
        AscendC::DataCopy(xLocal, xGlobal, dataSize);
        inQueue.EnQue(xLocal);
    }

    __aicore__ inline void Compute() {
        AscendC::LocalTensor<half> xLocal = inQueue.DeQue<half>();
        AscendC::LocalTensor<half> yLocal = outQueue.AllocTensor<half>();
        AscendC::Add(yLocal, xLocal, xLocal, dataSize);
        outQueue.EnQue<half>(yLocal);
        inQueue.FreeTensor(xLocal);
    }

    __aicore__ inline void CopyOut() {
        AscendC::LocalTensor<half> yLocal = outQueue.DeQue<half>();
        AscendC::DataCopy(yGlobal, yLocal, dataSize);
        outQueue.FreeTensor(yLocal);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue;
    AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue;
    AscendC::GlobalTensor<half> xGlobal, yGlobal;
    uint32_t dataSize;
};

AscendC SynchronizationControl InterfaceSummary

Overview

SynchronizationControlUsed for AI Core InternalAsynchronousParallelExecutesingleunitofbetweenofcoordinateadjust, distributeascoreinsideSynchronizationandcorebetweenSynchronization.

⚠️ relatedkeygeneralmiss: DMA Asynchronous

MTE2 (GM→UB) and MTE3 (UB→GM) isAsynchronousof: DataCopyPad returnreturntimeDataData Copystillnotcompleted.

mustmustin DataCopyPad ofafterthroughexceed EnQue/DeQue Synchronization, onlyabilitysafeallvisitaskData:

AllocTensor → DataCopyPad → EnQue(VECIN)    ← standardrememberData Copycompleted
                            DeQue(VECIN)     ← etcwaitData Copycompleted, onlyabilityCalculation
                            ... Calculation ...
                            EnQue(VECOUT)    ← standardrememberCalculationcompleted
                            DeQue(VECOUT)    ← etcwaitCalculationcompleted, onlyabilitytransferoutput
                            DataCopyPad → FreeTensor

pushrecommend EnQue/DeQue andnon PipeBarrier: EnQue/DeQue isprecisionconfirmofliveproduce-disappearcostSynchronization, PipeBarrier isthickparticledegreeallPipelinelineblockblock.

---

one, coreinsideSynchronization

1.1 Pipelinetypetype

1.2 multiplePipelineSynchronization (SetFlag/WaitFlag)

functionability: notsamePipelinelinebetweenofSynchronization, Used forDataaccordingdependScenarios.

ISASIInterface (notkeepcertifystrideVersionCompatibility) :

template <HardEvent event>
__aicore__ inline void SetFlag(int32_t eventID);

template <HardEvent event>
__aicore__ inline void WaitFlag(int32_t eventID);

TQueSyncInterface (keepcertifystrideVersionCompatibility) :

AscendC::TQueSync<PIPE_S, PIPE_MTE3> sync;
sync.SetFlag(0);
sync.WaitFlag(0);

HardEventtypetype: MTE2_V, V_MTE2, MTE3_V, V_MTE3, M_V, V_M, S_MTE3etc

Usageneedpoint:

• SetFlag/WaitFlagmustmustbecomepairoutputappear
• eventIDneedthroughexceedAllocEventID()orFetchEventID()obtainget
• rangescope: Atlastrainpractice0-3, otherother0-7

Examples:

dstLocal.SetValue(0, 0);
int32_t eventID = GetTPipePtr()->FetchEventID(AscendC::HardEvent::S_MTE3);
AscendC::SetFlag<AscendC::HardEvent::S_MTE3>(eventID);
AscendC::WaitFlag<AscendC::HardEvent::S_MTE3>(eventID);
AscendC::DataCopy(dstGlobal, dstLocal, dataSize);

1.3 singlePipelineSynchronization (PipeBarrier)

functionability: sameonePipelinelineInternalSynchronization, keepcertifypreviousorderfingercommandcompletedafterExecuteaftercontinuefingercommand.

reasontype:

template <pipe_t pipe>
__aicore__ inline void PipeBarrier()

Examples:

AscendC::Add(dst0Local, src0Local, src1Local, 512);
AscendC::PipeBarrier<PIPE_V>();  // keepcertifyAddcompleted
AscendC::Mul(dst1Local, dst0Local, src2Local, 512);

Notes: PIPE_SProhibitedadjustusePipeBarrier, ableleadissueHardwareerrorerror.

1.4 DataSynchronizationscreenobstacle (DataSyncBarrier)

functionability: blockblockaftercontinuefingercommandstraighttoallhaveMemoryvisitaskcompleted.

reasontype:

template <MemDsbT arg0>
__aicore__ inline void DataSyncBarrier()

parameternumber: ALL(allhaveMemory), DDR(GM), UB, SEQ(prekeep)

Examples:

AscendC::Mmad(...);
AscendC::DataSyncBarrier<MemDsbT::ALL>();
AscendC::Fixpipe(...);

---

two, corebetweenSynchronization

2.1 SyncAll (allcoreSynchronization)

functionability: allhavecoreSynchronization, etcwaitallhavecoreExecution Complete.

yieldingSynchronization:

template <bool isAIVOnly = true>
__aicore__ inline void SyncAll(const GlobalTensor<int32_t>& gmWorkspace,
                               const LocalTensor<int32_t>& ubWorkspace,
                               const int32_t usedCores = 0);

rigidSynchronization:

template <bool isAIVOnly = true>
__aicore__ inline void SyncAll();

emptybetweenRequirements: gmWorkspace ≥ Core count×32Bytes, ubWorkspace ≥ Core count×32Bytes

2.2 IBSet/IBWait (corebetweenSynchronization)

functionability: corebetweenSet/etcwaitSynchronizationstandardwill.

reasontype:

template <bool isAIVOnly = true>
__aicore__ inline void IBSet(const GlobalTensor<int32_t>& gmWorkspace,
                             const LocalTensor<int32_t>& ubWorkspace,
                             int32_t blockIdx, int32_t eventID);

template <bool isAIVOnly = true>
__aicore__ inline void IBWait(...);  // parameternumbersameabove

emptybetweenRequirements: gmWorkspace ≥ Core count×32×eventID_max + blockIdx_max×32 + 32

2.3 DeterminepropertyCalculationInterface

InitDetermineComputeWorkspace - InitializationtogethershareMemory:

__aicore__ inline void InitDetermineComputeWorkspace(
    GlobalTensor<int32_t>& gmWorkspace,
    LocalTensor<int32_t>& ubWorkspace);

WaitPreBlock - etcwaitpreviousonecountcore:

__aicore__ inline void WaitPreBlock(
    GlobalTensor<int32_t>& gmWorkspace,
    LocalTensor<int32_t>& ubWorkspace);

NotifyNextBlock - throughknowunderonecountcore:

__aicore__ inline void NotifyNextBlock(
    GlobalTensor<int32_t>& gmWorkspace,
    LocalTensor<int32_t>& ubWorkspace);

2.4 CrossCoreSetFlag/CrossCoreWaitFlag (distributeleavemodelformula)

functionability: surfacedirectiondistributeleavemodelformulaofcorebetweenSynchronization.

reasontype:

template <uint8_t modeId, pipe_t pipe>
__aicore__ inline void CrossCoreSetFlag(uint16_t flagId);

template <uint8_t modeId = 0, pipe_t pipe = PIPE_S>
__aicore__ inline void CrossCoreWaitFlag(uint16_t flagId);

modeId:

• 0: AI CorecorebetweenSynchronization
• 1: AIVcoreofbetweenSynchronization
• 2: AICandAIVofbetweenSynchronization

Examples:

// modelformula0: SynchronizationallhaveAIVcore
AscendC::CrossCoreSetFlag<0x0, PIPE_MTE3>(0x8);
AscendC::CrossCoreWaitFlag(0x8);

---

three, EventIDManagement

AllocEventID

applypleaseandoccupyuseEventID, needmatchmatchReleaseEventIDexplainrelease:

AscendC::TEventID eventID = GetTPipePtr()->AllocEventID<AscendC::HardEvent::V_S>();

FetchEventID

onlyobtaingetcanuseEventID, notoccupyuse:

AscendC::TEventID eventID = GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_S);

---

produceproductSupportSummary

AscendC VectorCalculation API Reference

API Calling Pattern

VectorCalculation API Provides Threeadjustusemethods:

1. Whole Tensor participationCalculation (runcomputecharacterweightload)

dstLocal = src0Local + src1Local;  // Add
dstLocal = src0Local < src1Local;  // Compare

2. Tensor previous n countDataCalculation

AscendC::Add(dstLocal, src0Local, src1Local, count);

3. Tensor highdimensionTilingCalculation

// connectcontinuemodelformula
AscendC::Add(dstLocal, src0Local, src1Local, mask, repeatTime, repeatParams);
// Iterativebitmodelformula
AscendC::Add(dstLocal, src0Local, src1Local, maskArray, repeatTime, repeatParams);

mask parameternumber

ControleachtimeiteraterepresentparticipationCalculationofunitelement:

// connectcontinuemodelformula
uint64_t mask = 128;  // handlemanageprevious128countunitelement

// Iterativebitmodelformula
uint64_t mask[2] = {UINT64_MAX, UINT64_MAX};  // handlemanageallpart128countunitelement

repeatParams parameternumber

BinaryRepeatParams (doublesourceoperateworknumber)

AscendC::BinaryRepeatParams {dstBlkStride, src0BlkStride, src1BlkStride,
                              dstRepStride, src0RepStride, src1RepStride};

UnaryRepeatParams (singlesourceoperateworknumber)

AscendC::UnaryRepeatParams {dstBlkStride, srcBlkStride, dstRepStride, srcRepStride};

CommonConfiguration: connectcontinueDatahandlemanage

• half: {1, 1, 1, 8, 8, 8}
• float: {1, 1, 8, 8} (UnaryRepeatParams)

Basecomputeart API

twounitruncompute

// Add, Sub, Mul, Div, Max, Min
AscendC::Add(dstLocal, src0Local, src1Local, count);
AscendC::Mul(dstLocal, src0Local, src1Local, mask, repeatTime, {1, 1, 1, 8, 8, 8});

// runcomputecharacterweightload
dstLocal = src0Local + src1Local;

Supporttypetype: half, int16_t, int32_t, float

oneunitruncompute

// Abs, Exp, Ln, Sqrt, Rsqrt, Reciprocal, Relu, LeakyRelu, Tanh
AscendC::Exp(dstLocal, srcLocal, count);

standardamountruncompute (optimizefirstUsage)

// Adds, Muls, Maxs, Mins — straightconnectpaireachcountunitelementdostandardamountoperatework
AscendC::Adds(dstLocal, srcLocal, scalarValue, count);
AscendC::Muls(dstLocal, srcLocal, scalarValue, count);

standardamountOptimization: travelreturnapproximatelyafterrequirespaireachtraveldecreaseremove/dividein order toonecountstandardamountvaluetime, optimizefirstuse Adds/Muls:

multipletravelwidebroadcast (BinaryRepeatParams)

pairmultipletravelDatadosameonecountdirectionamountofruncompute (ifeachtraveldecreaseremove max directionamount) :

// src1RepStride=0 use src1 ineachtime repeat timenotpreviousenter, Implementationwidebroadcast
uint64_t mask = alignedCols / (32 / sizeof(float));  // each repeat handlemanageofunitelementnumber
uint32_t repeatTime = rowCount;
AscendC::Sub(dst, src0, src1, mask, repeatTime,
             {1, 1, 1,                                    // blkStride
              alignedCols / 8, alignedCols / 8, 0});      // repStride: src1=0 widebroadcast

⚠️ repeatTime Limitations: repeatTime parameternumbertypetypeas uint8_t, mostlargevalue 255. exceedexceedneeddistributebatchhandlemanage:

while (remaining > 0) {
    uint8_t batch = static_cast<uint8_t>(std::min(remaining, (int64_t)255));
    AscendC::Sub(dst[offset], src0[offset], src1, mask, batch, params);
    offset += batch * alignedCols;
    remaining -= batch;
}

typetypeconvertexchange API (Cast)

AscendC::Cast(dstLocal, srcLocal, AscendC::RoundMode::CAST_RINT, count);

RoundMode give upinputmodelformula

Commonconvertexchangegroupmatch

mixmatchprecisiondegreemodelformula (FP16/BF16 upgradeprecisiondegreeCalculation)

returnapproximately/returnonetransformtypeOperator (Softmax, LayerNorm etc) requires FP32 middlebetweenprecisiondegreekeepcertifynumbervaluestabledefineproperty:

// Init Phase: amountexternaldistributematch FP32 workworkslowconflictregion
pipe.InitBuffer(calcBuf, alignedCols * sizeof(float));  // FP32 Calculationemptybetween

// Compute Phase: Iterativetravel Cast → FP32 Calculation → Cast return
LocalTensor<half> inLocal = inQueue.DeQue<half>();
LocalTensor<float> workLocal = calcBuf.Get<float>();
AscendC::Cast(workLocal, inLocal[rowIdx * alignedCols], RoundMode::CAST_NONE, rLength);
// ... FP32 returnapproximately/Calculation ...
AscendC::Cast(outLocal[rowIdx * alignedCols], workLocal, RoundMode::CAST_ROUND, rLength);

highdimensionTilingExamples

// half -> int32_t
uint64_t mask = 64;  // in order toint32_tasstandard
AscendC::Cast(dstLocal, srcLocal, AscendC::RoundMode::CAST_CEIL, mask, 8, {1, 1, 8, 4});

returnapproximatelyCalculation API

Level 2 returnapproximately (singletravel/anymeaninglengthdegree)

// ReduceSum / ReduceMax / ReduceMin
// tmpBuffer typetypemustmustand T mutualsame, notabilityis uint8_t
AscendC::ReduceSum(dstLocal, srcLocal, sharedTmpBuffer, count);
AscendC::ReduceMax(dstLocal, srcLocal, sharedTmpBuffer, count);

• count = havevalidunitelementnumber (rLength) , notispairalignafteroflengthdegree
• tmpBuffer typetypemustmustis LocalTensor<T> (andsourcemutualsametypetype)
• dst notabilityand tmpBuffer fingerdirectionsameoneblockMemory

tmpBuffer largesmallCalculation

int elementsPerBlock = 32 / sizeof(T);      // half:16, float:8
int elementsPerRepeat = 256 / sizeof(T);    // half:128, float:64
int firstMaxRepeat = (count + elementsPerRepeat - 1) / elementsPerRepeat;
int tmpBufferSize = ((firstMaxRepeat + elementsPerBlock - 1) / elementsPerBlock) * elementsPerBlock;
// distributematch: pipe.InitBuffer(tmpBuf, tmpBufferSize * sizeof(T));

Pattern returnapproximately (batchamountmultipletravel 2D)

pairalignDataofbatchamounttravelreturnapproximately, Performanceupdateoptimize:

// Pattern::Reduce::AR — returnapproximatelymostafteronedimension (eachtravelreturnapproximatelyasonecountstandardamount) 
// srcShape = {rows, alignedCols}, alignedCols mustmust 32B pairalign
AscendC::ReduceMax(dstLocal, srcLocal, sharedTmpBuffer,
                   srcShape, AscendC::Pattern::Reduce::AR, srcInnerPad);

tmpBuffer largesmall: Usage GetReduceMaxMaxMinTmpSize / GetReduceSumTmpSize Calculation:

uint32_t tmpSize = AscendC::GetReduceMaxMaxMinTmpSize<T>(srcShape);
pipe.InitBuffer(tmpBuf, tmpSize);

⚠️ Pattern returnapproximatelyRequirements alignedCols is 32B pairalignof, nonpairalignDataUsage Level 2 Iterativetravelreturnapproximately.

WholeReduceSum / BlockReduceSum

Hardwarefingercommand, PerformanceupdateoptimizebutLimitationsupdatemultiple:

AscendC::WholeReduceSum(dstLocal, srcLocal, count);

comparecompareandselectselect API

Compare

// runcomputecharacterweightload
dstLocal = src0Local < src1Local;

// functionnumberadjustuse
AscendC::Compare(dstLocal, src0Local, src1Local, AscendC::CMPMODE::LT, count);

CMPMODE: LT(<), GT(>), GE(>=), LE(<=), EQ(==), NE(!=)

Output: uint8_t typetype, according to bit bitMemoryResult

⚠️ 256 charactersectionpairalignConstraints: Compare API RequirementsparticipationcomparecompareofDataregiondomainis 256 charactersectionofadjustnumbertimes. nonpairaligntimerequires padding:

// notfoot 256B ofpartdistributefillfill ±inf / FLT_MAX, confirmkeep padding regionnotshadowloudResult
uint32_t alignedCount = ((count * sizeof(T) + 255) / 256) * (256 / sizeof(T));
AscendC::Duplicate(src[count], paddingValue, alignedCount - count);  // padding
AscendC::Compare(dst, src0, src1, CMPMODE::LT, alignedCount);

Select

// modelformula0: twocounttensorselectget (selMaskhavebitnumberLimitations) 
AscendC::Select(dstLocal, maskLocal, src0Local, src1Local,
                AscendC::SELMODE::VSEL_CMPMASK_SPR, count);

// modelformula1: tensorandscalarselectget
AscendC::Select(dstLocal, maskLocal, src0Local, scalarValue,
                AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, count);

// modelformula2: twocounttensorselectget (selMaskconnectcontinuedisappearconsume) 
AscendC::Select(dstLocal, maskLocal, src0Local, src1Local,
                AscendC::SELMODE::VSEL_TENSOR_TENSOR_MODE, count);

selMask Rules: bit bitas1select src0, as0select src1

Datafillfill API (Duplicate)

AscendC::Duplicate(dstLocal, scalarValue, count);

// highdimensionTiling
AscendC::Duplicate(dstLocal, scalarValue, mask, repeatTime, dstBlkStride, dstRepStride);

complexmatchCalculation API

// FusedMulAdd: dst = src0 * src1 + src2
AscendC::FusedMulAdd(dstLocal, src0Local, src1Local, src2Local, count);

// FusedMulAddRelu: dst = Relu(src0 * src1 + src2)
AscendC::FusedMulAddRelu(dstLocal, src0Local, src1Local, src2Local, count);

// Axpy: dst = a * x + y
AscendC::Axpy(dstLocal, aLocal, xLocal, yLocal, count);

CommonCodemodelformula

unitelementlevelruncompute

__aicore__ inline void Compute()
{
    LocalTensor<half> src0 = inQueueX.DeQue<half>();
    LocalTensor<half> src1 = inQueueY.DeQue<half>();
    LocalTensor<half> dst = outQueueZ.AllocTensor<half>();

    AscendC::Add(dst, src0, src1, tileLength);

    outQueueZ.EnQue(dst);
    inQueueX.FreeTensor(src0);
    inQueueY.FreeTensor(src1);
}

upgradeprecisiondegreeCalculation (FP16 -> FP32)

__aicore__ inline void Compute()
{
    LocalTensor<half> src0 = inQueueX.DeQue<half>();
    LocalTensor<half> src1 = inQueueY.DeQue<half>();
    LocalTensor<half> dst = outQueueZ.AllocTensor<half>();

    // complexuseMemoryperformtypetypeconvertexchange
    LocalTensor<float> src0Fp32 = src0.ReinterpretCast<float>();
    LocalTensor<float> src1Fp32 = src1.ReinterpretCast<float>();
    LocalTensor<float> dstFp32 = dst.ReinterpretCast<float>();

    AscendC::Cast(src0Fp32, src0, AscendC::RoundMode::CAST_NONE, tileLength);
    AscendC::Cast(src1Fp32, src1, AscendC::RoundMode::CAST_NONE, tileLength);
    AscendC::Add(dstFp32, src0Fp32, src1Fp32, tileLength);
    AscendC::Cast(dst, dstFp32, AscendC::RoundMode::CAST_NONE, tileLength);

    outQueueZ.EnQue(dst);
    inQueueX.FreeTensor(src0);
    inQueueY.FreeTensor(src1);
}

Prerequisitesselectselect

__aicore__ inline void Compute()
{
    LocalTensor<float> src0 = inQueueX.DeQue<float>();
    LocalTensor<float> src1 = inQueueY.DeQue<float>();
    LocalTensor<uint8_t> cmpResult = tmpQueue.AllocTensor<uint8_t>();
    LocalTensor<float> dst = outQueueZ.AllocTensor<float>();

    // comparecompare
    AscendC::Compare(cmpResult, src0, src1, AscendC::CMPMODE::LT, tileLength);
    // selectselect
    AscendC::Select(dst, cmpResult, src0, src1,
                    AscendC::SELMODE::VSEL_CMPMASK_SPR, tileLength);

    outQueueZ.EnQue(dst);
    inQueueX.FreeTensor(src0);
    inQueueY.FreeTensor(src1);
    tmpQueue.FreeTensor(cmpResult);
}

throughuseConstraints

• regionaddresspairalign: LocalTensor risebeginregionaddressneed 32 charactersectionpairalign
• Datatypetypeoneconsistent: sourceoperateworknumberanditemofoperateworknumbertypetypeneedoneconsistent (Cast divideexternal)
• TPosition: Support VECIN/VECCALC/VECOUT
• repeatTime ≤ 255: UsagehighdimensionTilingmodelformulatime, repeatTime as uint8_t, transferinput >255 ablequietsilentcutjudgeas 0 guideconsistenterrorerrorResult. needin host SideLimitationsor kernel Sidedistributebatch
• dst ≠ tmpBuffer: ReduceMax/ReduceSum of dst notabilityand tmpBuffer issameoneblockMemory
• Prohibited std:: Math Functions: Kernel middleProhibited std::min/max/abs/sqrt/exp etc, Usage AscendC directionamount API orthreeunitruncomputecharactersubstituterepresent

Phase 2: Data CopyOptimization — DetailedReference

2.1 Single Transfer Size >= 16 KB

Bandwidth UtilizationfollowSingle Transfer Sizeincreaselargeandliftupgrade. actualtestExperience: SingleData Copy >= 16 KB time, UB↔HBM twocountmethoddirectionaveragecanreachtoconnectnearpeakvalueofbandwidthwidth. lowinthisvaluetimeBandwidth Utilizationdisplayfamousunderdowngrade.

setcalculate Tiling strategystrategytimeshouldconfirmkeepeachtime DataCopy Data Copyarrivefew 16 KB.

2.2 GM regionaddress 512B pairalign

in Atlas A2 trainpracticesystemcolumn / Atlas 800I A2 pushmanageproduceproductabove, GM regionaddress 512B pairaligncancompare 32B pairalignobtainobtainmosthigh 30% ofbandwidthwidthliftupgrade (mostdifferenceScenariosunderofdifferencedistance) .

distributematch GM Tensor orCalculationpartialmoveamounttime, shouldconfirmkeeprisebegincharactersectionregionaddressas 512 ofadjustnumbertimes.

2.3 Usage stride parameternumberrepresentsubstitute for loopring

Usage DataCopyParams (blockCount / blockLen / srcStride / dstStride) willbetweenseparate Data CopyDescriptionasoneitem DMA fingercommandunderissue, andnonuse for loopringIterativetraveladjustuse DataCopy.

reverseexample — for loopring, eachtimeonlyData Copy 2 KB:

constexpr int32_t copyWidth = 2 * 1024 / sizeof(float);
constexpr int32_t imgWidth  = 16 * 1024 / sizeof(float);
constexpr int32_t imgHeight = 16;
// 16 timealoneestablishof 2KB Data Copy, Bandwidth Utilizationextremelow
for (int i = 0; i < imgHeight; i++) {
    DataCopy(tensorIn[i * copyWidth], tensorGM[i * imgWidth], copyWidth);
}

positiveexample — singleitem DMA Descriptioncharacter, onetimeData Copy 32 KB:

constexpr int32_t copyWidth = 2 * 1024 / sizeof(float);
constexpr int32_t imgWidth  = 16 * 1024 / sizeof(float);
constexpr int32_t imgHeight = 16;
DataCopyParams copyParams;
copyParams.blockCount = imgHeight;                     // 16 travel
copyParams.blockLen   = copyWidth / 8;                 // singlebit: 32B DataBlock
copyParams.srcStride  = (imgWidth - copyWidth) / 8;    // src travelbetweenbetweenseparate
copyParams.dstStride  = 0;                             // dst connectcontinuecomposeinput
DataCopy(tensorGM, tensorIn, copyParams);

stride methodsunderissueoneitem DMA fingercommand, HardwareselfmaincompletedallpartData Copy, canfilldistributeutilizeusebandwidthwidth. for loopringmethodsunderissue 16 itemsmall DMA fingercommand, eachitemofbetweenalsohave Scalar openconsume.

Phase 4: MemoryOptimization — DetailedReference

Memory HierarchyOverview

4.1 UB Buffer mergematch

connectcontinuemultipletime Vector runcomputetime, willmiddlebetweenResultkeepkeepin UB above, notvia GM towardsreturn. n timeconnectcontinueruncomputeof GM Data Copytimenumberfrom 2n downgradeas 2.

reverseexample — eachtimeruncomputeallvia GM towardsreturn (Exp + Abs need 4 time GM Data Copy) :

class KernelSample {
    __aicore__ inline void Process() {
        CopyIn();       // GM → UB
        Compute();      // Exp
        CopyOut();      // UB → GM
        CopyIn1();      // GM → UB (weightnewreadreturn Exp Result) 
        Compute1();     // Abs
        CopyOut1();     // UB → GM
    }
};

positiveexample — in UB insidelinkformulaCalculation (only 2 time GM Data Copy) :

class KernelSample {
    __aicore__ inline void Compute() {
        LocalTensor<float> src0Local = inQueueSrc0.DeQue<float>();
        LocalTensor<float> dstLocal  = outQueueDst.AllocTensor<float>();
        Exp(dstLocal, src0Local, 1024);
        Abs(dstLocal, dstLocal, 1024);   // reasonregionoperatework, keepin UB
        outQueueDst.EnQue<float>(dstLocal);
        inQueueSrc0.FreeTensor(src0Local);
    }
    __aicore__ inline void Process() {
        CopyIn();       // GM → UB (onetime) 
        Compute();      // Exp + Abs mergematch
        CopyOut();      // UB → GM (onetime) 
    }
};

4.2 L0C accumulateaddrulematrixmultiply

A1*B1 + A2*B2 + ... Scenariosunder, utilizeuse Mmad ofinsidebuildaccumulateaddfunctionabilitywillpartdistributeResultkeepkeepin CO1 (L0C) middle. avoidavoideachtimerulematrixmultiplyResultall CO1→GM→UB againdo Add.

reverseexample — Iterativetimetransferoutputafterin UB requestand:

void Process() {
    Compute();       // Mmad → CO1
    CopyOut();       // CO1 → workspace (GM) 
    CopyIn1();       // workspace → UB

    Compute1();      // Mmad → CO1
    CopyOut1();      // CO1 → workspace (GM) 
    CopyIn2();       // workspace → UB

    Compute2();      // Add(result1, result2) in UB
    CopyOut2();      // UB → GM
}

positiveexample — in L0C middlereasonregionaccumulateadd:

void Compute() {
    MmadParams mmadParams;
    mmadParams.m = m;  mmadParams.n = n;  mmadParams.k = k;
    Mmad(c1Local, a2Local_1, b2Local_1, mmadParams);
    mmadParams.cmatrixInitVal = false;
    Mmad(c1Local, a2Local_2, b2Local_2, mmadParams);  // in CO1 reasonregionaccumulateadd
}
// mostafteronetime CopyOut: CO1 → GM

4.3 smallrulematrixlengthstation L1

when L1 nomethodsametimecontentcontainleftrightrulematrix (ifleftrulematrix 992K, rightrulematrix 16K, L1 contentamount 512K) time, willcomparesmallrulematrixonetimeLoadafteroftenstation L1, onlyloopringData Copycomparelargerulematrix.

reverseexample — eachtimeiteraterepresentallweightnewLoadtwocountrulematrix:

void Process() {
    for (uint32_t i = 0; i < 2; i++) {
        CopyInA1(i);      // Loadleftrulematrixcutslice
        SplitA();
        for (uint32_t j = 0; j < 2; j++) {
            CopyInB1(j);  // eachtimeallweightnewLoadrightrulematrix
            SplitB();
            Compute(i, j);
        }
    }
    CopyOut();
}

positiveexample — rightrulematrixonetimeLoad, onlyloopringData Copyleftrulematrix:

void Process() {
    CopyInB1();          // rightrulematrixonetimeallloadinput L1
    SplitB();            // L1 → L0B
    for (uint32_t i = 0; i < 2; i++) {
        CopyInA1(i);     // loopringLoadleftrulematrixcutslice
        SplitA();
        for (uint32_t j = 0; j < 2; j++) {
            Compute(i, j);  // rightrulematrixalreadyin L0B
        }
    }
    CopyOut();
}

2 countleftrulematrixcutslicetime: Data Copytimenumberfrom 4+4=8 downgradeas 1+2=3.

4.4 BT Buffer Store bias (distributeleaveArchitecture)

will bias keepinput BT Buffer (C2) , in Mmad middleonestepmergematch bias addmethod, avoidavoid CO1→GM→UB→Add→GM ofredundantlengthPath.

reverseexample — in UB middlesinglealonedo bias Add:

TQue<QuePosition::VECIN, 1> inQueueBias;
// Mmad after: CO1 → workspace(GM) → UB
// bias: GM → UB
// Add(matmul_result, bias) in UB → GM

positiveexample — throughexceed BT Buffer mergematch:

TQue<QuePosition::C1, 1> inQueueC1;    // L1
TQue<QuePosition::C2, 1> outQueueC2;   // BT Buffer

void SplitBias() {
    LocalTensor<float> bias1Local = inQueueC1.DeQue<float>();
    LocalTensor<float> bias2Local = outQueueC2.AllocTensor<float>();
    // L1 → BT
    DataCopy(bias2Local, bias1Local, {1, (uint16_t)(n * sizeof(float) / 64), 0, 0});
    outQueueC2.EnQue<float>(bias2Local);
    inQueueC1.FreeTensor(bias1Local);
}

void Compute() {
    LocalTensor<float> bias2Local = outQueueC2.DeQue<float>();
    MmadParams mmadParams;
    mmadParams.m = m;  mmadParams.n = n;  mmadParams.k = k;
    mmadParams.cmatrixInitVal = false;
    Mmad(c1Local, a2Local, b2Local, bias2Local, mmadParams);  // mergematch bias
    outQueueC2.FreeTensor(bias2Local);
}

4.5 FP Buffer StoreQuantizationparameternumber (distributeleaveArchitecture)

willQuantizationparameternumberkeepinput FP Buffer (C2PIPE2GM) , throughexceed Fixpipe intransferoutputPathabovefollowpathQuantization. avoidavoid CO1→GM→UB→QuantizationCalculation→GM ofredundantlengthPath.

reverseexample — in UB middlesinglealonedoQuantization:

TQue<QuePosition::VECIN, 1> inQueueDeq;    // Quantizationparameternumberin UB
// CO1 → workspace → UB
// Quantizationparameternumber: GM → UB
// Cast + Mul + Cast in UB
// Result → GM

positiveexample — throughexceed FP Buffer mergematch:

TQue<QuePosition::C1, 1>        inQueueDeq1;   // L1
TQue<QuePosition::C2PIPE2GM, 1> inQueueDeq;    // FP Buffer

void SplitDeq() {
    LocalTensor<uint64_t> deq1Local = inQueueDeq1.DeQue<uint64_t>();
    LocalTensor<uint64_t> deqLocal  = inQueueDeq.AllocTensor<uint64_t>();
    // L1 → FP Buffer
    DataCopy(deqLocal, deq1Local, {1, (uint16_t)(cSize * sizeof(uint64_t) / 128), 0, 0});
    inQueueDeq.EnQue<uint64_t>(deqLocal);
    inQueueDeq1.FreeTensor(deq1Local);
}

void CopyOut() {
    LocalTensor<float>    c1Local  = outQueueCO1.DeQue<float>();
    LocalTensor<uint64_t> deqLocal = inQueueDeq.DeQue<uint64_t>();
    SetFixpipeNz2ndFlag(1, 0, 0);
    DataCopyCO12DstParams params;
    params.nSize    = n;
    params.mSize    = m;
    params.srcStride = m;
    params.dstStride = n;
    params.quantPre = QuantMode_t::VQF322B8_PRE;
    params.nz2ndEn  = true;
    DataCopy(cGM, c1Local, params);   // transferoutputtimefollowpathQuantization
    outQueueCO1.FreeTensor(c1Local);
}

Phase 5: PipelineOptimization — DetailedReference

5.1 CopyIn / Compute / CopyOut Paradigm

willOperatorDivide IntoThree-level PipelineTask, Usage TQue performlevelbetweenSynchronization. notsamePhasereflectshoottoaloneestablishof Hardwarefingercommandteamcolumn (MTE2/MTE3 Data Copy, V Vector, M rulematrix) , canParallelExecute.

CopyIn  → AllocTensor + DataCopy(GM→Local) + EnQue     [MTE2 teamcolumn]
Compute → DeQue + Vector/Cube runcompute + EnQue              [V / M teamcolumn]
CopyOut → DeQue + DataCopy(Local→GM) + FreeTensor       [MTE3 teamcolumn]

BaseFramework:

TPipe pipe;
TQue<VecIn, 1> queIn;
TQue<VecOut, 1> queOut;

pipe.InitBuffer(queIn, 2, 1024);   // double buffer

for (int i = 0; i < tileCount; i++) {
    // CopyIn
    auto tensor = queIn.AllocTensor<half>();
    DataCopy(tensor, gm, len);
    queIn.EnQue(tensor);

    // Compute
    auto tensorIn = queIn.DeQue<half>();
    auto tensorOut = queOut.AllocTensor<half>();
    Abs(tensorOut, tensorIn, 1024);
    queIn.FreeTensor(tensorIn);
    queOut.EnQue(tensorOut);

    // CopyOut
    auto result = queOut.DeQue<half>();
    DataCopy(gmOut, result, 1024);
    queOut.FreeTensor(result);
}

sameoneDatacutsliceinside, CopyIn → Compute → CopyOut mustmuststringtravel. butnotsamecutslicecanweightstack: Compute handlemanagecutslice N time, CopyIn cantransferinputcutslice N+1, CopyOut cantransferoutputcutslice N−1.

5.2 Double Buffer

InitBuffer of buffer countnumbersetas 2, useData CopyandCalculationweightstackExecute.

reverseexample — notuseability double buffer (Vector utilizeuserateapproximately 33%) :

pipe.InitBuffer(inQueueSrc0, 1, sizeSrc0 * sizeof(half));   // single buffer
pipe.InitBuffer(inQueueSrc1, 1, sizeSrc1 * sizeof(half));
pipe.InitBuffer(outQueueDst, 1, sizeDst0 * sizeof(half));

for (uint32_t index = 0; index < round * 2; ++index) {
    CopyIn(index);    // MTE2 busy, Vector idle
    Compute();        // Vector busy, MTE idle
    CopyOut(index);   // MTE3 busy, Vector idle
}

positiveexample — useability double buffer (underone tile of CopyIn andwhenprevious tile of Compute weightstack) :

pipe.InitBuffer(inQueueSrc0, 2, sizeSrc0 * sizeof(half));   // double buffer
pipe.InitBuffer(inQueueSrc1, 2, sizeSrc1 * sizeof(half));
pipe.InitBuffer(outQueueDst, 2, sizeDst0 * sizeof(half));

for (uint32_t index = 0; index < round; ++index) {
    CopyIn(index);    // canandpreviousonetime CopyOut weightstack
    Compute();        // canandunderonetime CopyIn weightstack
    CopyOut(index);   // canandunderonetime Compute weightstack
}

Notesmatteritem:

• Memoryopenconsumefliptimes (eachcountteamcolumndistributematch 2 block buffer) .
• loopringtimenumbermust >= 2 onlyabilityobtainadvantageous.
• whenCalculationtimebetweenfarlargeinData Copytimebetweentime, Data Copyalreadybehiddenhidden, double buffer receiveadvantageoushavelimit.
• whenDataamountverysmall, onetimeimmediatecancompletedallpartCalculationtime, noneed double buffer.

5.3 Asynchronous Iterate (MIX modelformula, AIC+AIV)

Matmul MIX Scenariosunder, Iterate / IterateAll ablein AIV (Vector core) and AIC (Cube core) ofbetweenissuepresentSynchronizationdisappearinformation. SynchronizationmodelformulaControldisappearinformationfrequencyrate:

• Iterate<true>() (Synchronization) : eachtimeiteraterepresentissueoneitemdisappearinformation——openconsumelarge.
• Iterate<false>() (Asynchronous) : onlychapteronetimeissuedisappearinformation, aftercontinueiteraterepresentnoneed AIC/AIV Synchronization.

Synchronizationmodelformula — eachtimeiteraterepresentallhavedisappearinformationopenconsume:

AIV: send_msg → wait → send_msg → wait → send_msg → wait
AIC:           exec →           exec →           exec

Asynchronousmodelformula — onlyfirsttimeissuedisappearinformation:

AIV: send_msg → continue → continue → continue
AIC:           exec     → exec     → exec

CodeExamples:

TQueBind<TPosition::CO2, TPosition::VECIN>    qVecIn;
TQueBind<TPosition::VECIN, TPosition::VECOUT> qVecOut;

mm.SetTensorA(gmA);
mm.SetTensorB(gmB);
mm.SetWorkspace(workspace, singleCoreM * singleCoreN * sizeof(float));

while (mm.template Iterate<false>()) {    // Asynchronousmodelformula
    auto cInUB = qVecIn.AllocTensor<float>();
    mm.GetTensorC(cInUB);
    qVecIn.EnQue(cInUB);
    cInUB = qVecIn.Deque<float>();
    auto cOutUB = qVecOut.AllocTensor<float>();
    Muls(cOutUB, cInUB, scalar, baseM * baseN);
    qVecIn.FreeTensor(cInUB);
    // ... aftercontinuehandlemanage
}

MIX ScenariosundersilentrecognizeUsageAsynchronousmodelformula. onlyinrequiresstrictformatofIterativetimeiteraterepresentsequenceorderkeepcertify (guardstopregionaddressstampstep) timeonlyFallbackasSynchronizationmodelformula.

Phase 6: Scalar — DetailedReference

6.1 Vector transformOptimization (Eliminate Scalar Loop)

Problem Pattern: Usage for loopringIterativeunitelementoperatework (if one-hot compilecodeofloopringassignvalue)

oftenseeOptimizationMethod:

• use Duplicate + Single SetValue substituterepresentloopringassignvalue
• use Exp, Log, Mul, Muls etcdirectionQuantization API substituterepresentIterativeunitelementoperatework

Examples (one-hot compilecodeOptimization) :

reverseexample — for loopring, eachtimethroughexceedSetValueperformstandardamountsubstituteexchangeoperatework: :

for (size_t classIdx = 0; classIdx < numClass_; ++classIdx) {
    int64_t weight = 0;
    if (labelIdx == classIdx) {
        weight = 1;
    }
    labelOneHotLocal_.SetValue(classIdx, weight);
}

positiveexample — UsageDuplicatebatchamountsubstituteexchange:

Duplicate(labelOneHotLocal_, float(0.0), numClass_);
// Synchronizationkeepcertify Duplicate completedafteragain SetValue
PipeBarrier<PIPE_V>();
TEventID eventIdVToS = GetTPipePtr()->FetchEventID(HardEvent::V_S);
SetFlag<HardEvent::V_S>(eventIdVToS);
WaitFlag<HardEvent::V_S>(eventIdVToS);
labelOneHotLocal_.SetValue(labelIdx, float(1.0));

Optimizationreasonmanage:

• Duplicate isoneitemdirectionQuantizationfingercommand, HardwareParallelfillfillWhole tensor, compareloopringIterativecount SetValue rapidnumbertentimes
• loopringmethodsundereachcount iteration allhavestandardamountjudgejudgeandstandardamountassignvalueopenconsume, Scalar fingercommandoccupycompareextremehigh
• Optimizationafteronlyneed 1 item Vector fingercommand + 1 itemstandardamountassignvalue, largerangedowngradelow Scalar occupycompare

6.2 Vector transformOptimization (Eliminate Scalar Loop)

Usage DataCopyParams (blockCount / blockLen / srcStride / dstStride) willbetweenseparate Data CopyDescriptionasoneitem DMA fingercommandunderissue, andnonuse for loopringIterativetraveladjustuse DataCopy.

reverseexample — for loopring, eachtimeonlyData Copy 2 KB:

constexpr int32_t copyWidth = 2 * 1024 / sizeof(float);
constexpr int32_t imgWidth  = 16 * 1024 / sizeof(float);
constexpr int32_t imgHeight = 16;
// 16 timealoneestablishof 2KB Data Copy, Bandwidth Utilizationextremelow
for (int i = 0; i < imgHeight; i++) {
    DataCopy(tensorIn[i * copyWidth], tensorGM[i * imgWidth], copyWidth);
}

positiveexample — singleitem DMA Descriptioncharacter, onetimeData Copy 32 KB:

constexpr int32_t copyWidth = 2 * 1024 / sizeof(float);
constexpr int32_t imgWidth  = 16 * 1024 / sizeof(float);
constexpr int32_t imgHeight = 16;
DataCopyParams copyParams;
copyParams.blockCount = imgHeight;                     // 16 travel
copyParams.blockLen   = copyWidth / 8;                 // singlebit: 32B DataBlock
copyParams.srcStride  = (imgWidth - copyWidth) / 8;    // src travelbetweenbetweenseparate
copyParams.dstStride  = 0;                             // dst connectcontinuecomposeinput
DataCopy(tensorGM, tensorIn, copyParams);

Optimizationreasonmanage:

• stride methodsunderissueoneitem DMA fingercommand, HardwareselfmaincompletedallpartData Copy, canfilldistributeutilizeusebandwidthwidth.
• for loopringmethodsunderissue 16 itemsmall DMA fingercommand, eachitemofbetweenalsohave Scalar openconsume.

Phase 1: Tiling Optimization — DetailedReference

1.1 Multi-core Tiling

throughexceed context->SetBlockDim(BLOCK_DIM) SetOperatorUsageofCore count.

blockDim aslogiclogiccoregeneralmiss, getvaluerangescope [1, 65535]. asfilldistributeutilizeuseHardwareResource, onegeneralsetas objectmanageCore countorotheradjustnumbertimes. AIC/AIV Core countdistributecategorythroughexceed GetCoreNumAic() and GetCoreNumAiv() obtainget.

1.2 L2Cache Tiling

when outputinputDataamount + OutputDataamount > L2Cache contentamount (if 192 MB) time, willDataaccording to L2Cache largesmalletcdistributeasmultipleblock, allhavecorecoordinatesamehandlemanagesameoneblockafteragaincutexchangeunderoneblock. thissampleweightcomplexReadtimecancommandmiddle L2Cache (~7 TB/s) , avoidavoidfrequencycomplexvisitask HBM (~1.6 TB/s) .

reverseexample — notuseability L2Cache Tiling, eachcountcoreoftwocount tile mutualmutualcrowdoccupy L2Cache:

constexpr int32_t TOTAL_LENGTH = 384 * 1024 * 1024 / sizeof(half);
constexpr int32_t USE_CORE_NUM = 20;
constexpr int32_t TILE_NUM = 2;
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM;

class KernelSample {
public:
    __aicore__ inline void Init(GM_ADDR x) {
        xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
        pipe.InitBuffer(inQueueX, 1, BLOCK_LENGTH * sizeof(half));
    }
    __aicore__ inline void Process() {
        constexpr int32_t loopCount = 2;
        for (int32_t i = 0; i < loopCount; i++) {
            for (int32_t j = 0; j < TILE_NUM; j++) {
                CopyIn(j);    // eachcountcorereadtwocount tile, L2Cache bereversecomplexwasheliminate
                Compute();
                CopyOut(j);
            }
        }
    }
};

positiveexample — useability L2Cache Tiling, externallayerloopringaccording to L2Cache distributeblock, allhavecorecoordinatesamehandlemanage:

constexpr int32_t TOTAL_LENGTH = 384 * 1024 * 1024 / sizeof(half);
constexpr int32_t TILE_NUM = 2;
constexpr int32_t USE_CORE_NUM = 20;
constexpr int32_t TILE_LENGTH = TOTAL_LENGTH / TILE_NUM;
constexpr int32_t BLOCK_LENGTH = TILE_LENGTH / USE_CORE_NUM;

class KernelSample {
public:
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, int32_t index) {
        xGm.SetGlobalBuffer(
            (__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx() + index * TILE_LENGTH,
            BLOCK_LENGTH);
    }
    __aicore__ inline void Process() {
        constexpr int32_t loopCount = 2;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn();       // eachcountcoreonlyreadselfselfofcutslice, chaptertwotimereadcommandmiddle L2Cache
            Compute();
            CopyOut();
        }
    }
};

extern "C" __global__ __aicore__ void simple_kernel(
    __gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
{
    AscendC::KernelAdd op;
    for (int32_t i = 0; i < TILE_NUM; i++) {
        op.Init(srcGm, dstGm, i);
        op.Process();
    }
}

1.3 corebetweennegativeloadaveragebalance

L2Cache Tilingafter, ifeachtimeCalculationallneedblocknumbernotabilitybeCore countadjustdivide, rulepartdistributecoreablemultipledistributematchtailblock.

askproblem: core 1–5 eachtime pass multiplecomputeonecountblock, beginfinalmostaftercompleted, core 6–20 emptyetc.

solveresolveMethod: innotsame pass betweenexchangesubstitutedistributematchtailblock. exampleif 2 count pass × 25 block / 20 core, pass 1 oftailblockdistributematchgivecore 1–5, pass 2 oftailblockdistributematchgivecore 6–10, Globalcomeseecore 1–10 eachcompute 3 block, core 11–20 eachcompute 2 block, reachtoGlobalnegativeloadaveragebalance.

Troubleshooting

1. Compilation Issues

Issue: Header file not found

Symptom: fatal error: ascendc/kernel.h: No such file or directory

Solution:

# Set include path
export ASCEND_INCLUDE_PATH=/usr/local/Ascend/ascend-toolkit/latest/compiler/include
export CPLUS_INCLUDE_PATH=$ASCEND_INCLUDE_PATH:$CPLUS_INCLUDE_PATH

# Or use cmake with proper include paths

Issue: Linker error undefined reference

Symptom: undefined reference to 'ascendc::xxx'

Solution:

# Link with ascendc library
export ASCEND_LIB_PATH=/usr/local/Ascend/ascend-toolkit/latest compiler/lib64
export LD_LIBRARY_PATH=$ASCEND_LIB_PATH:$LD_LIBRARY_PATH

# Or update CMakeLists.txt
target_link_libraries(op_test ascendc)

Issue: CMake not finding toolchain

Symptom: Could not find toolchain

Solution:

# Set toolchain file
cmake -DCMAKE_TOOLCHAIN_FILE=../toolchain.cmake ..

# Or verify toolchain path
cat toolchain.cmake | grep CMAKE_C_COMPILER

2. Runtime Issues

Issue: ACL initialization failed

Symptom: aclError: 1

Solution:

# Set ACL config path
export ASCEND_CONFIG_PATH=/usr/local/Ascend/ascend-toolkit/latest/
export LD_LIBRARY_PATH=$ASCEND_CONFIG_PATH/acllib/lib64:$LD_LIBRARY_PATH

# Initialize before running
aclInit(nullptr);

Issue: Memory allocation failed

Symptom: Failed to allocate Tensor

Solution:

// Check available memory
// Reduce buffer sizes
// Free unused buffers before allocation

Issue: Kernel launch failed

Symptom: Kernel launch error

Solution:

# Check NPU status
npu-smi info

# Verify operator compiled for correct chip
# Check soc-version matches hardware

3. Profiling Issues

Issue: OPPROF directory not created

Root Cause: Insufficient permissions or wrong path

Solution:

# Use absolute path
msprof op --output=/tmp/opprof ./execute_op

# Check directory permissions
chmod 777 /tmp/opprof

Issue: Profiling data incomplete

Root Cause: Application crashes or too short execution

Solution:

# Increase iteration count
# Add warmup loop
# Extend execution time

4. Optimization Issues

Issue: Performance not improved

Possible Causes: 1. Optimization target incorrect 2. Change too small to measure 3. Bottleneck elsewhere

Solution:

# Re-profile to confirm bottleneck
msprof op --aic-metrics=Roofline --output=./opprof ./execute_op

# Check PipeUtilization forPipelineline efficiency
# Check Memory forbandwidthwidth bottlenecks

Issue: Accuracy degraded

Root Cause: Wrong optimization logic

Solution:

# Revert changes
cp operator_dir_backup_*/op_kernel/*.cpp operator_dir/op_kernel/

# Rebuild and retest
bash build.sh
./execute_op

Issue: Anti-pattern violation not detected

Root Cause: Not checking all rules

Solution:

# Review anti-pattern checklist:
# - FP16/BF16 for complex math
# - No right-value in EXEC_KERNEL_CMD
# - No GM<->UB DataCopy
# - No reuse after ReduceSum/ReduceMax
# - No std::min/max/sqrt/exp in kernel

5. Debug Commands

# Check environment
echo $ASCEND_TOOLKIT_HOME
echo $LD_LIBRARY_PATH

# Verify headers
ls /usr/local/Ascend/ascend-toolkit/latest/compiler/include/ascendc/

# Test basic operator
./execute_op

# Check output format
cat op_output.json

Quick Diagnostic

#!/bin/bash
echo "=== Environment ==="
cmake --version 2>/dev/null || echo "cmake: NOT FOUND"
aarch64-linux-gnu-g++ --version 2>/dev/null | head -1 || echo "compiler: NOT FOUND"

echo "=== Source ==="
ls -la op_kernel/*.cpp 2>/dev/null || echo "No kernel source"

echo "=== Build ==="
bash build.sh 2>&1 | tail -5

echo "=== Run ==="
./execute_op 2>&1 | tail -10

Verification Methods

Prerequisite Verification

1. Verify Development Environment

# Check CANN installation
cat /usr/local/Ascend/ascend-toolkit/latest/version.ini

# Check AscendC headers
ls -la /usr/local/Ascend/ascend-toolkit/latest/compiler/include/ascendc/

# Verify ACL availability
ls -la /usr/local/Ascend/ascend-toolkit/latest/

2. Verify Operator Source

# Check operator directory structure
ls -la operator_dir/
# Expected: op_host/, op_kernel/, CMakeLists.txt

# Check source files
ls -la operator_dir/op_kernel/*.cpp

3. Verify Build Tools

# Check cmake
cmake --version

# Check compiler
aarch64-linux-gnu-g++ --version

# Check build script
cat operator_dir/build.sh

Functional Verification

Phase 1: Investigation Verification

# Read operator design document
cat operator_dir/design_doc.md

# Read source code
cat operator_dir/op_kernel/add.cpp

# Generate investigation report
# Should list optimization points by phase

Phase 2: Baseline Verification

# Backup operator directory
cp -r operator_dir operator_dir_backup_$(date +%Y%m%d)

# Run baseline profiling
cd operator_dir
msprof op --output=./baseline_opprof ./execute_op

# Verify baseline report
ls -la *_baseline_report.md

Phase 3: Optimization Verification

# Verify reference loaded
# Check ascendc-api references

# Apply code modifications
vi op_kernel/add.cpp
vi op_host/add.cpp

# Rebuild
bash build.sh

# Verify compilation
# Should show "Build success" or no errors

Phase 4: Accuracy Verification

# Run accuracy test
cd operator_dir
./execute_op

# Check output
# Should show "pass" or "accuracy OK"
# Should not show errors

# Compare with baseline
echo "Accuracy verification: PASS"

Phase 5: Performance Verification

# Collect post-optimization data
cd operator_dir
msprof op --output=./optim_opprof ./execute_op

# Generate comparison
# Should show before/after metrics

# Calculate speedup
# (baseline - optimized) / baseline * 100%

End-to-End Verification Script

#!/bin/bash
set -e

echo "=== 1. Verify Environment ==="
cmake --version
aarch64-linux-gnu-g++ --version

echo "=== 2. Verify Operator Source ==="
ls -la operator_dir/op_kernel/*.cpp

echo "=== 3. Backup Original ==="
cp -r operator_dir operator_dir_backup_$(date +%Y%m%d%H%M%S)

echo "=== 4. Build Baseline ==="
cd operator_dir && bash build.sh

echo "=== 5. Run Baseline Profiling ==="
msprof op --output=./baseline ./execute_op

echo "=== 6. Apply Optimizations ==="
# Code modifications here

echo "=== 7. Rebuild ==="
bash build.sh

echo "=== 8. Verify Accuracy ==="
./execute_op

echo "=== 9. Run Post-Optim Profiling ==="
msprof op --output=./optimized ./execute_op

echo "=== All verifications passed ==="

Verification Checklist