Star1280
1391
代码
代码
Issues77
Pull Requests192
Actions
项目讨论
Wiki
分析
Star1280
1391
已合并
slice,pad,pad_v3,mirror_pad修改tilingdata注册 #1595
huangzhiyuan创建于 3月11日
slice,pad,pad_v3,mirror_pad修改tilingdata注册 #1595
已合并
huangzhiyuan创建于 3月11日
从h1234515/ops-mathslice_faketilingdata_change合入到cann/ops-mathmaster
共 6 个文件变更+9-10
Mconversion/mirror_pad/op_kernel/mirror_pad_apt.cpp+2-1
@@ -51,7 +51,7 @@ extern "C" __global__ __aicore__ void mirror_pad(
51 GM_ADDR x, GM_ADDR paddings, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)51 GM_ADDR x, GM_ADDR paddings, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
52{52{
53 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);53 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
54- REGISTER_TILING_DEFAULT(SliceFakeTilingData);54+ REGISTER_NONE_TILING;
55 55 
56 if (TILING_KEY_IS(REFLECT_SIMT_BRANCH)) { // 2100056 if (TILING_KEY_IS(REFLECT_SIMT_BRANCH)) { // 21000
57 PadV3::LaunchKernelPadMirrorSimt<DTYPE_X, REFLECT_SIMT_BRANCH>(x, paddings, y, tiling);57 PadV3::LaunchKernelPadMirrorSimt<DTYPE_X, REFLECT_SIMT_BRANCH>(x, paddings, y, tiling);
@@ -119,6 +119,7 @@ extern "C" __global__ __aicore__ void mirror_pad(
119 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);119 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);
120 PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);120 PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);
121 } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {121 } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {
122+ GET_TILING_DATA_WITH_STRUCT(SliceTilingData, tilingData, tiling);
122 // 空tenseor处理123 // 空tenseor处理
123 } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {124 } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {
124 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);125 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
Mconversion/pad/op_kernel/pad_apt.cpp+2-1
@@ -48,7 +48,7 @@ extern "C" __global__ __aicore__ void pad(GM_ADDR x, GM_ADDR paddings, GM_ADDR y
48 }48 }
49 SetSysWorkspace(workspace);49 SetSysWorkspace(workspace);
50 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);50 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
51- REGISTER_TILING_DEFAULT(SliceFakeTilingData);51+ REGISTER_NONE_TILING;
52 if (TILING_KEY_IS(CONSTANT_CUT_LAST_DIM_BRANCH)) { // 3000052 if (TILING_KEY_IS(CONSTANT_CUT_LAST_DIM_BRANCH)) { // 30000
53 PadV3::LaunchKernelPadWithHugeWidth<DTYPE_X>(x, paddings, y, tiling);53 PadV3::LaunchKernelPadWithHugeWidth<DTYPE_X>(x, paddings, y, tiling);
54 } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM2)) { // 3002154 } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM2)) { // 30021
@@ -91,6 +91,7 @@ extern "C" __global__ __aicore__ void pad(GM_ADDR x, GM_ADDR paddings, GM_ADDR y
91 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);91 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);
92 PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);92 PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);
93 } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {93 } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {
94+ GET_TILING_DATA_WITH_STRUCT(SliceTilingData, tilingData, tiling);
94 // 空tenseor处理95 // 空tenseor处理
95 } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {96 } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {
96 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);97 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
Mconversion/pad_v3/op_kernel/pad_v3_apt.cpp+2-1
@@ -89,7 +89,7 @@ extern "C" __global__ __aicore__ void pad_v3(
89 GM_ADDR x, GM_ADDR paddings, GM_ADDR constValues, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)89 GM_ADDR x, GM_ADDR paddings, GM_ADDR constValues, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
90{90{
91 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);91 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
92- REGISTER_TILING_DEFAULT(SliceFakeTilingData);92+ REGISTER_NONE_TILING;
93 if (TILING_KEY_IS(CONSTANT_CUT_LAST_DIM_BRANCH)) { // 3000093 if (TILING_KEY_IS(CONSTANT_CUT_LAST_DIM_BRANCH)) { // 30000
94 PadV3::LaunchKernelPadWithHugeWidth<DTYPE_X>(x, paddings, y, tiling, constValues);94 PadV3::LaunchKernelPadWithHugeWidth<DTYPE_X>(x, paddings, y, tiling, constValues);
95 } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM2)) { // 3002195 } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM2)) { // 30021
@@ -227,6 +227,7 @@ extern "C" __global__ __aicore__ void pad_v3(
227 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);227 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);
228 PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);228 PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);
229 } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {229 } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {
230+ GET_TILING_DATA_WITH_STRUCT(SliceTilingData, tilingData, tiling);
230 // 空tenseor处理231 // 空tenseor处理
231 } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {232 } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {
232 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);233 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
Mconversion/slice/op_host/arch35/slice_tiling_arch35.cpp+1-1
@@ -358,7 +358,7 @@ void SliceTiling::FillSliceTilingData101()
358 size_t tilingDataSize = sizeof(SliceMoveAlignLastDimTilingData);358 size_t tilingDataSize = sizeof(SliceMoveAlignLastDimTilingData);
359 OP_LOGD(tilingContext_->GetNodeName(), "Entering FillTilingData101.");359 OP_LOGD(tilingContext_->GetNodeName(), "Entering FillTilingData101.");
360 FillSliceBaseTilingData(sliceMoveAlignLastDimTilingData_.sliceBaseTilingData);360 FillSliceBaseTilingData(sliceMoveAlignLastDimTilingData_.sliceBaseTilingData);
361- auto tilingData = tilingContext_->GetTilingData<SliceMoveAlignTilingData>();361+ auto tilingData = tilingContext_->GetTilingData<SliceMoveAlignLastDimTilingData>();
362 errno_t ret = memcpy_s(362 errno_t ret = memcpy_s(
363 tilingData, tilingDataSize, reinterpret_cast<void*>(&sliceMoveAlignLastDimTilingData_), tilingDataSize);363 tilingData, tilingDataSize, reinterpret_cast<void*>(&sliceMoveAlignLastDimTilingData_), tilingDataSize);
364 if (ret != EOK) {364 if (ret != EOK) {
Mconversion/slice/op_kernel/arch35/slice_struct.h+0-5
@@ -21,11 +21,6 @@ constexpr int64_t MAX_AXIS_NUM_FOR_STRIDESLICE = 8;
21constexpr int64_t MAX_NDDMA_UB_SPLIT_AXIS_NUM = 5;21constexpr int64_t MAX_NDDMA_UB_SPLIT_AXIS_NUM = 5;
22constexpr int64_t MAX_SIMT_UB_SPLIT_AXIS_NUM = 8;22constexpr int64_t MAX_SIMT_UB_SPLIT_AXIS_NUM = 8;
23constexpr int64_t NUMBER_TWO = 2;23constexpr int64_t NUMBER_TWO = 2;
24-constexpr int64_t MAX_TILINGDATA_BYTES = 4096;
25- 
26-struct SliceFakeTilingData {
27- char scalarData[MAX_TILINGDATA_BYTES];
28-};
29 24 
30struct SliceBaseTilingData {25struct SliceBaseTilingData {
31 int8_t isBeginConst;26 int8_t isBeginConst;
Mconversion/slice/op_kernel/slice_apt.cpp+2-1
@@ -31,7 +31,7 @@ extern "C" __global__ __aicore__ void slice(
31 GM_ADDR x, GM_ADDR offsets, GM_ADDR size, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)31 GM_ADDR x, GM_ADDR offsets, GM_ADDR size, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
32{32{
33 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);33 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
34- REGISTER_TILING_DEFAULT(SliceFakeTilingData);34+ REGISTER_NONE_TILING;
35 TPipe pipe;35 TPipe pipe;
36 if (TILING_KEY_IS(SLICE_KEY_MOVE_ALIGN)) {36 if (TILING_KEY_IS(SLICE_KEY_MOVE_ALIGN)) {
37 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignTilingData, tilingData, tiling);37 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignTilingData, tilingData, tiling);
@@ -49,6 +49,7 @@ extern "C" __global__ __aicore__ void slice(
49 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);49 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);
50 SliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);50 SliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);
51 } else if (TILING_KEY_IS(SLICE_KEY_SIMT)) {51 } else if (TILING_KEY_IS(SLICE_KEY_SIMT)) {
52+ GET_TILING_DATA_WITH_STRUCT(SliceTilingData, tilingData, tiling);
52 // 空tenseor处理53 // 空tenseor处理
53 } else if (TILING_KEY_IS(SLICE_KEY_MOVE_ALIGN_GATHER)) {54 } else if (TILING_KEY_IS(SLICE_KEY_MOVE_ALIGN_GATHER)) {
54 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);55 GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);