CheckLocalMemoryIA
功能说明
监视设定范围内的UB读写行为,如果监视到有设定范围的读写行为则会出现EXCEPTION报错,未监视到设定范围的读写行为则不会报错。
函数原型
1
|
__aicore__ inline void CheckLocalMemoryIA(const CheckLocalMemoryIAParam& checkParams) |
参数说明
参数名称 |
输入/输出 |
含义 |
||
---|---|---|---|---|
checkParams |
输入 |
矩阵乘相关参数,类型为CheckLocalMemoryIAParam,结构体具体定义为:
参数说明请参考表2。 |
参数名称 |
含义 |
---|---|
enableBit |
配置的异常寄存器,取值范围:enableBit∈[0,3],默认为0。
|
startAddr |
Check的起始地址,32B对齐,取值范围:startAddr∈[0, 65535],默认值为0。比如,可通过LocalTensor.GetPhyAddr()/32来获取startAddr。 |
endAddr |
Check的结束地址,32B对齐,取值范围:startAddr∈[0, 65535] 。默认值为0。 |
isScalarRead |
Check标量读访问。
|
isScalarWrite |
Check标量写访问。
|
isVectorRead |
Check矢量读访问。
|
isVectorWrite |
Check矢量写访问。
|
isMteRead |
Check Mte读访问。
|
isMteWrite |
Check Mte写访问。
|
isEnable |
是否使能enableBit参数配置的异常寄存器。
|
约束说明
- startAddr/endAddr的单位是32B,监控范围不包含startAddr,包含endAddr,即(startAddr,endAddr]。
- 每次调用完该接口需要进行复位(配置isEnable为false进行复位);
- 操作数地址偏移对齐要求请参见通用约束。
支持的型号
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas推理系列产品(Ascend 310P处理器)AI Core
Atlas 200/500 A2推理产品
调用示例
该示例check矢量写访问是否在设定的(startAddr, endAddr]范围内。当前示例check到矢量写在设定的范围内,结果会报错(ACL_ERROR_RT_VECTOR_CORE_EXECPTION)。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 |
#include "kernel_operator.h" class KernelAdd { public: __aicore__ inline KernelAdd() {} __aicore__ inline void Init(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm) { src0Global.SetGlobalBuffer((__gm__ half*)src0Gm); src1Global.SetGlobalBuffer((__gm__ half*)src1Gm); dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm); pipe.InitBuffer(inQueueSrc0, 1, 512 * sizeof(half)); pipe.InitBuffer(inQueueSrc1, 1, 512 * sizeof(half)); pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(half)); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensor<half> src0Local = inQueueSrc0.AllocTensor<half>(); AscendC::LocalTensor<half> src1Local = inQueueSrc1.AllocTensor<half>(); AscendC::DataCopy(src0Local, src0Global, 512); AscendC::DataCopy(src1Local, src1Global, 512); inQueueSrc0.EnQue(src0Local); inQueueSrc1.EnQue(src1Local); } __aicore__ inline void Compute() { AscendC::LocalTensor<half> src0Local = inQueueSrc0.DeQue<half>(); AscendC::LocalTensor<half> src1Local = inQueueSrc1.DeQue<half>(); AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>(); AscendC::CheckLocalMemoryIA({ 0, (uint32_t)(dstLocal.GetPhyAddr() / 32), (uint32_t)((dstLocal.GetPhyAddr() + 512 * sizeof(half)) / 32), false, false, false, true, false, false, true }); AscendC::Add(dstLocal, src0Local, src1Local, 512); outQueueDst.EnQue<half>(dstLocal); inQueueSrc0.FreeTensor(src0Local); inQueueSrc1.FreeTensor(src1Local); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>(); AscendC::DataCopy(dstGlobal, dstLocal, 512); outQueueDst.FreeTensor(dstLocal); } private: AscendC::TPipe pipe; AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc0, inQueueSrc1; AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst; AscendC::GlobalTensor<half> src0Global, src1Global, dstGlobal; }; extern "C" __global__ __aicore__ void add_simple_kernel(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm) { KernelAdd op; op.Init(src0Gm, src1Gm, dstGm); op.Process(); } |
父主题: 矩阵计算