Star915
1833
代码
代码
Issues230
Pull Requests324
Actions
项目讨论
Wiki
分析
Star915
1833
已合并
matmul类算子kernel入口添加InitSocState初始化寄存器状态 #6445
AlbertYoung192创建于 6月24日
matmul类算子kernel入口添加InitSocState初始化寄存器状态 #6445
已合并
AlbertYoung192创建于 6月24日
从AlbertYoung192/ops-nn_shangfen9.1.0合入到cann/ops-nn9.1.0
共 6 个文件变更+6-0
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3.cpp+1-0
@@ -166,6 +166,7 @@ template <
166__global__ __aicore__ void batch_mat_mul_v3(166__global__ __aicore__ void batch_mat_mul_v3(
167 GM_ADDR aGM, GM_ADDR bGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR cGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)167 GM_ADDR aGM, GM_ADDR bGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR cGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)
168{168{
169+ AscendC::InitSocState();
169 __gm__ uint8_t* user = GetUserWorkspace(workspaceGM);170 __gm__ uint8_t* user = GetUserWorkspace(workspaceGM);
170 171 
171 constexpr bool aTran = (BATCH_A_TRANS == 1);172 constexpr bool aTran = (BATCH_A_TRANS == 1);
Mmatmul/fused_mat_mul/op_kernel/fused_mat_mul.cpp+1-0
@@ -181,6 +181,7 @@ template <
181__global__ __aicore__ void fused_mat_mul(181__global__ __aicore__ void fused_mat_mul(
182 GM_ADDR x1GM, GM_ADDR x2GM, GM_ADDR biasGM, GM_ADDR x3GM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)182 GM_ADDR x1GM, GM_ADDR x2GM, GM_ADDR biasGM, GM_ADDR x3GM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)
183{183{
184+ AscendC::InitSocState();
184 __gm__ uint8_t* user = AscendC::GetUserWorkspace(workspaceGM);185 __gm__ uint8_t* user = AscendC::GetUserWorkspace(workspaceGM);
185 186 
186 constexpr bool aTran = (TRANS_MODEL == 1 || TRANS_MODEL == 3);187 constexpr bool aTran = (TRANS_MODEL == 1 || TRANS_MODEL == 3);
Mmatmul/mat_mul_v3/op_kernel/arch35/mat_mul_v3.cpp+1-0
@@ -93,6 +93,7 @@ template <
93__global__ __aicore__ void mat_mul_v3(93__global__ __aicore__ void mat_mul_v3(
94 GM_ADDR aGM, GM_ADDR bGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR cGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)94 GM_ADDR aGM, GM_ADDR bGM, GM_ADDR biasGM, GM_ADDR offsetWGM, GM_ADDR cGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)
95{95{
96+ AscendC::InitSocState();
96 constexpr bool aTran = (A_TRANS == 1);97 constexpr bool aTran = (A_TRANS == 1);
97 constexpr bool bTran = (B_TRANS == 1);98 constexpr bool bTran = (B_TRANS == 1);
98 99 
Mmatmul/rotate_quant/op_kernel/arch35/rotate_quant.cpp+1-0
@@ -25,6 +25,7 @@ using RotateQuantAptOpt::RotateQuantAptTilingData;
25extern "C" __global__ __aicore__ void rotate_quant(25extern "C" __global__ __aicore__ void rotate_quant(
26 GM_ADDR x, GM_ADDR rot, GM_ADDR alpha, GM_ADDR y, GM_ADDR scale, GM_ADDR workSpace, GM_ADDR tiling)26 GM_ADDR x, GM_ADDR rot, GM_ADDR alpha, GM_ADDR y, GM_ADDR scale, GM_ADDR workSpace, GM_ADDR tiling)
27{27{
28+ AscendC::InitSocState();
28 REGISTER_TILING_DEFAULT(RotateQuantAptTilingData);29 REGISTER_TILING_DEFAULT(RotateQuantAptTilingData);
29 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);30 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
30 GET_TILING_DATA_WITH_STRUCT(RotateQuantAptTilingData, tilingData, tiling);31 GET_TILING_DATA_WITH_STRUCT(RotateQuantAptTilingData, tilingData, tiling);
Mmatmul/transpose_batch_mat_mul/op_kernel/arch35/transpose_batch_mat_mul.cpp+1-0
@@ -54,6 +54,7 @@ template <int8_t PERM_X1, int8_t PERM_X2, int8_t BATCH_SPLIT>
54__global__ __aicore__ void transpose_batch_mat_mul(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR biasGM, GM_ADDR scalesGM,54__global__ __aicore__ void transpose_batch_mat_mul(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR biasGM, GM_ADDR scalesGM,
55 GM_ADDR cGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)55 GM_ADDR cGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)
56{56{
57+ AscendC::InitSocState();
57 __gm__ uint8_t* user = GetUserWorkspace(workspaceGM);58 __gm__ uint8_t* user = GetUserWorkspace(workspaceGM);
58 59 
59 constexpr bool aTran = false;60 constexpr bool aTran = false;
Mquant/flat_quant/op_kernel/flat_quant_apt.cpp+1-0
@@ -28,6 +28,7 @@ extern "C" __global__ __aicore__ void flat_quant(
28 GM_ADDR x, GM_ADDR kronecker_p1, GM_ADDR kronecker_p2, GM_ADDR out, GM_ADDR quant_scale, GM_ADDR workspace,28 GM_ADDR x, GM_ADDR kronecker_p1, GM_ADDR kronecker_p2, GM_ADDR out, GM_ADDR quant_scale, GM_ADDR workspace,
29 GM_ADDR tiling)29 GM_ADDR tiling)
30{30{
31+ AscendC::InitSocState();
31 GET_TILING_DATA(tilingData, tiling);32 GET_TILING_DATA(tilingData, tiling);
32 const FlatQuantTilingData* __restrict tiling_data = &tilingData;33 const FlatQuantTilingData* __restrict tiling_data = &tilingData;
33 const TCubeTiling* __restrict mmTilingR = &(tiling_data->matmulTilingR);34 const TCubeTiling* __restrict mmTilingR = &(tiling_data->matmulTilingR);