Star486
625
代码
代码
Issues43
Pull Requests196
Actions
项目讨论
Wiki
分析
Star486
625
已合并
add aiv print #2667
yevolin2023创建于 22 天前
add aiv print #2667
已合并
yevolin2023创建于 22 天前
从yevolin2023/hcclmaster合入到cann/hcclmaster
共 26 个文件变更+129-80
MCMakeLists.txt+6-0
@@ -34,12 +34,18 @@ else()
34 set(ENABLE_DEVICE OFF)34 set(ENABLE_DEVICE OFF)
35endif()35endif()
36 36 
37+# AIV算子调测阶段可配置为1以打开AscendC打印功能
38+if(NOT DEFINED ASCENDC_DUMP)
39+ set(ASCENDC_DUMP "0")
40+endif()
41+ 
37message(STATUS "ASCEND_CANN_PACKAGE_PATH=${ASCEND_CANN_PACKAGE_PATH}")42message(STATUS "ASCEND_CANN_PACKAGE_PATH=${ASCEND_CANN_PACKAGE_PATH}")
38message(STATUS "CANN_3RD_LIB_PATH=${CANN_3RD_LIB_PATH}")43message(STATUS "CANN_3RD_LIB_PATH=${CANN_3RD_LIB_PATH}")
39message(STATUS "PRODUCT_SIDE=${PRODUCT_SIDE}")44message(STATUS "PRODUCT_SIDE=${PRODUCT_SIDE}")
40message(STATUS "BUILD_OPEN_PROJECT=${BUILD_OPEN_PROJECT}")45message(STATUS "BUILD_OPEN_PROJECT=${BUILD_OPEN_PROJECT}")
41message(STATUS "ENABLE_BUILD_DEVICE=${ENABLE_BUILD_DEVICE}")46message(STATUS "ENABLE_BUILD_DEVICE=${ENABLE_BUILD_DEVICE}")
42message(STATUS "ENABLE_BUILD_AARCH=${ENABLE_BUILD_AARCH}")47message(STATUS "ENABLE_BUILD_AARCH=${ENABLE_BUILD_AARCH}")
48+message(STATUS "ASCENDC_DUMP=${ASCENDC_DUMP}")
43 49 
44set(HCCL_DIR "${CMAKE_CURRENT_SOURCE_DIR}")50set(HCCL_DIR "${CMAKE_CURRENT_SOURCE_DIR}")
45set(HCC_TOOLCHAIN_DIR "${ASCEND_CANN_PACKAGE_PATH}/toolkit/toolchain/hcc")51set(HCC_TOOLCHAIN_DIR "${ASCEND_CANN_PACKAGE_PATH}/toolkit/toolchain/hcc")
Msrc/ops/all_gather/template/aiv/kernel/CMakeLists.txt+24-22
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_all_gather_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_gather_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_gather_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,7 +42,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int8_v2 ALL
41add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2.o42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2.o
42 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2.o43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2.o
43 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=044 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
44- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec45+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
45 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false46 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
46 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__47 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
47 -DHCCL_DTYPE_INT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int848 -DHCCL_DTYPE_INT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int8
@@ -62,7 +63,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int16_v2 ALL
62add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2.o63add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2.o
63 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2.o64 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2.o
64 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=065 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
65- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec66+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
66 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false67 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
67 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__68 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
68 -DHCCL_DTYPE_INT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int1669 -DHCCL_DTYPE_INT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int16
@@ -83,7 +84,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int32_v2 ALL
83add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2.o84add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2.o
84 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2.o85 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2.o
85 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=086 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
86- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec87+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
87 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false88 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
88 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__89 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
89 -DHCCL_DTYPE_INT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int3290 -DHCCL_DTYPE_INT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int32
@@ -104,7 +105,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_float16_v2 ALL
104add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2.o105add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2.o
105 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2.o106 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2.o
106 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0107 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
107- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec108+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
108 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false109 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
109 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__110 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
110 -DHCCL_DTYPE_FP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float16111 -DHCCL_DTYPE_FP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float16
@@ -125,7 +126,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_float32_v2 ALL
125add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2.o126add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2.o
126 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2.o127 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2.o
127 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0128 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
128- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec129+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
129 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false130 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
130 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__131 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
131 -DHCCL_DTYPE_FP32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float32132 -DHCCL_DTYPE_FP32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float32
@@ -146,7 +147,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_bfloat16_v2 ALL
146add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2.o147add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2.o
147 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2.o148 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2.o
148 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0149 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
149- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec150+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
150 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false151 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
151 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__152 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
152 -DHCCL_DTYPE_BFP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_bfloat16153 -DHCCL_DTYPE_BFP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_bfloat16
@@ -167,7 +168,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint8_v2 ALL
167add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2.o168add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2.o
168 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2.o169 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2.o
169 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0170 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
170- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec171+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
171 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false172 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
172 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__173 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
173 -DHCCL_DTYPE_UINT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint8174 -DHCCL_DTYPE_UINT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint8
@@ -188,7 +189,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint16_v2 ALL
188add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2.o189add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2.o
189 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2.o190 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2.o
190 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0191 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
191- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec192+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
192 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false193 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
193 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__194 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
194 -DHCCL_DTYPE_UINT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint16195 -DHCCL_DTYPE_UINT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint16
@@ -209,7 +210,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint32_v2 ALL
209add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2.o210add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2.o
210 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2.o211 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2.o
211 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0212 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
212- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec213+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
213 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false214 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
214 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__215 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
215 -DHCCL_DTYPE_UINT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint32216 -DHCCL_DTYPE_UINT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint32
@@ -230,7 +231,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int64_v2 ALL
230add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2.o231add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2.o
231 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2.o232 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2.o
232 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0233 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
233- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec234+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
234 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false235 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
235 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__236 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
236 -DHCCL_DTYPE_INT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int64237 -DHCCL_DTYPE_INT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int64
@@ -251,7 +252,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint64_v2 ALL
251add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2.o252add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2.o
252 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2.o253 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2.o
253 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0254 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
254- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec255+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
255 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false256 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
256 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__257 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
257 -DHCCL_DTYPE_UINT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint64258 -DHCCL_DTYPE_UINT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint64
@@ -338,6 +339,7 @@ add_custom_target(hccl_aiv_all_gather_op_910_96 ALL
338 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_96.o)339 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_96.o)
339add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_96.o340add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_96.o
340 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_gather_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec341 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_gather_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_gather_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
342+ -DASCENDC_DUMP=${ASCENDC_DUMP}
341 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000343 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
342 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/344 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
343 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw345 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
@@ -363,7 +365,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int8_v2_910_96 ALL
363add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2_910_96.o365add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2_910_96.o
364 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2_910_96.o366 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int8_v2_910_96.o
365 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0367 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
366- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec368+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
367 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false369 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
368 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__370 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
369 -DHCCL_DTYPE_INT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int8371 -DHCCL_DTYPE_INT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int8
@@ -391,7 +393,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int16_v2_910_96 ALL
391add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2_910_96.o393add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2_910_96.o
392 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2_910_96.o394 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int16_v2_910_96.o
393 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0395 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
394- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec396+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
395 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false397 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
396 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__398 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
397 -DHCCL_DTYPE_INT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int16399 -DHCCL_DTYPE_INT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int16
@@ -419,7 +421,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int32_v2_910_96 ALL
419add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2_910_96.o421add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2_910_96.o
420 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2_910_96.o422 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int32_v2_910_96.o
421 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0423 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
422- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec424+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
423 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false425 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
424 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__426 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
425 -DHCCL_DTYPE_INT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int32427 -DHCCL_DTYPE_INT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int32
@@ -447,7 +449,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_float16_v2_910_96 ALL
447add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2_910_96.o449add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2_910_96.o
448 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2_910_96.o450 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float16_v2_910_96.o
449 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0451 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
450- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec452+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
451 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false453 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
452 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__454 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
453 -DHCCL_DTYPE_FP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float16455 -DHCCL_DTYPE_FP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float16
@@ -475,7 +477,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_float32_v2_910_96 ALL
475add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2_910_96.o477add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2_910_96.o
476 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2_910_96.o478 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_float32_v2_910_96.o
477 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0479 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
478- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec480+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
479 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false481 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
480 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__482 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
481 -DHCCL_DTYPE_FP32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float32483 -DHCCL_DTYPE_FP32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_float32
@@ -503,7 +505,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_bfloat16_v2_910_96 ALL
503add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2_910_96.o505add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2_910_96.o
504 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2_910_96.o506 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_bfloat16_v2_910_96.o
505 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0507 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
506- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec508+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
507 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false509 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
508 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__510 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
509 -DHCCL_DTYPE_BFP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_bfloat16511 -DHCCL_DTYPE_BFP16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_bfloat16
@@ -531,7 +533,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint8_v2_910_96 ALL
531add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2_910_96.o533add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2_910_96.o
532 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2_910_96.o534 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint8_v2_910_96.o
533 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0535 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
534- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec536+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
535 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false537 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
536 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__538 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
537 -DHCCL_DTYPE_UINT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint8539 -DHCCL_DTYPE_UINT8 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint8
@@ -559,7 +561,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint16_v2_910_96 ALL
559add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2_910_96.o561add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2_910_96.o
560 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2_910_96.o562 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint16_v2_910_96.o
561 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0563 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
562- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec564+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
563 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false565 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
564 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__566 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
565 -DHCCL_DTYPE_UINT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint16567 -DHCCL_DTYPE_UINT16 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint16
@@ -587,7 +589,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint32_v2_910_96 ALL
587add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2_910_96.o589add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2_910_96.o
588 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2_910_96.o590 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint32_v2_910_96.o
589 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0591 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
590- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec592+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
591 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false593 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
592 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__594 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
593 -DHCCL_DTYPE_UINT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint32595 -DHCCL_DTYPE_UINT32 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint32
@@ -615,7 +617,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_int64_v2_910_96 ALL
615add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2_910_96.o617add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2_910_96.o
616 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2_910_96.o618 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_int64_v2_910_96.o
617 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0619 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
618- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec620+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
619 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false621 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
620 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__622 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
621 -DHCCL_DTYPE_INT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int64623 -DHCCL_DTYPE_INT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_int64
@@ -643,7 +645,7 @@ add_custom_target(hccl_ag_superkernel_mesh_1d_uint64_v2_910_96 ALL
643add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2_910_96.o645add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2_910_96.o
644 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2_910_96.o646 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ag_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ag_superkernel_mesh_1d_uint64_v2_910_96.o
645 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0647 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
646- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec648+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
647 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false649 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
648 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__650 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
649 -DHCCL_DTYPE_UINT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint64651 -DHCCL_DTYPE_UINT64 -Dsk_allgather_mesh_1d=sk_allgather_mesh_1d_uint64
Msrc/ops/all_gather/template/aiv/kernel/aiv_ag_mesh_1d_superkernel.h+1-0
@@ -13,6 +13,7 @@
13 13 
14extern "C" __aicore__ void sk_allgather_mesh_1d(SUPERKERNEL_LITE_ARGS_DEF)14extern "C" __aicore__ void sk_allgather_mesh_1d(SUPERKERNEL_LITE_ARGS_DEF)
15{15{
16+ AIV_INFO_HINT;
16 SUPERKERNEL_LITE_ARGS_EXTRACT;17 SUPERKERNEL_LITE_ARGS_EXTRACT;
17 return sk_ag_mesh_1d(SUPERKERNEL_ARGS_CALL);18 return sk_ag_mesh_1d(SUPERKERNEL_ARGS_CALL);
18}19}
Msrc/ops/all_gather/template/aiv/kernel/aiv_all_gather_op.h+1-0
@@ -19,6 +19,7 @@ using namespace AscendC;
19#define AIV_ALLGATHER_KERNEL_BATCH_DEF(type) \19#define AIV_ALLGATHER_KERNEL_BATCH_DEF(type) \
20 extern "C" __global__ __aicore__ void aiv_all_gather_##type(KERNEL_ARGS_DEF) \20 extern "C" __global__ __aicore__ void aiv_all_gather_##type(KERNEL_ARGS_DEF) \
21 { \21 { \
22+ AIV_INFO_HINT; \
22 return AivAllGatherV2Mesh1D<type>(KERNEL_ARGS_CALL); \23 return AivAllGatherV2Mesh1D<type>(KERNEL_ARGS_CALL); \
23 } \24 } \
24 EXPORT_AIV_META_INFO(aiv_all_gather_##type)25 EXPORT_AIV_META_INFO(aiv_all_gather_##type)
Msrc/ops/all_reduce/template/aiv/kernel/CMakeLists.txt+26-24
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_all_reduce_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,7 +42,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_int8_v2 ALL
41add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2.o42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2.o
42 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2.o43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2.o
43 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=044 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
44- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec45+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
45 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false46 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
46 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__47 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
47 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int848 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int8
@@ -62,7 +63,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_int16_v2 ALL
62add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2.o63add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2.o
63 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2.o64 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2.o
64 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=065 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
65- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec66+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
66 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false67 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
67 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__68 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
68 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int1669 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int16
@@ -83,7 +84,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_int32_v2 ALL
83add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2.o84add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2.o
84 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2.o85 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2.o
85 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=086 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
86- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec87+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
87 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false88 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
88 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__89 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
89 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int3290 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int32
@@ -104,7 +105,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_float16_v2 ALL
104add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2.o105add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2.o
105 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2.o106 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2.o
106 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0107 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
107- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec108+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
108 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false109 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
109 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__110 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
110 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float16111 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float16
@@ -125,7 +126,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_float32_v2 ALL
125add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2.o126add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2.o
126 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2.o127 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2.o
127 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0128 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
128- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec129+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
129 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false130 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
130 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__131 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
131 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float32132 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float32
@@ -146,7 +147,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2 ALL
146add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2.o147add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2.o
147 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2.o148 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2.o
148 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0149 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
149- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec150+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
150 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false151 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
151 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__152 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
152 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_bfloat16153 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_bfloat16
@@ -167,7 +168,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_int8_v2 ALL
167add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2.o168add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2.o
168 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2.o169 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2.o
169 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0170 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
170- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec171+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
171 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false172 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
172 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__173 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
173 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int8174 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int8
@@ -188,7 +189,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_int16_v2 ALL
188add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2.o189add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2.o
189 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2.o190 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2.o
190 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0191 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
191- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec192+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
192 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false193 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
193 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__194 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
194 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int16195 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int16
@@ -209,7 +210,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_int32_v2 ALL
209add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2.o210add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2.o
210 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2.o211 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2.o
211 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0212 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
212- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec213+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
213 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false214 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
214 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__215 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
215 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int32216 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int32
@@ -230,7 +231,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_float16_v2 ALL
230add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2.o231add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2.o
231 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2.o232 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2.o
232 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0233 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
233- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec234+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
234 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false235 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
235 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__236 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
236 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float16237 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float16
@@ -251,7 +252,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_float32_v2 ALL
251add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2.o252add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2.o
252 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2.o253 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2.o
253 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0254 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
254- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec255+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
255 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false256 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
256 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__257 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
257 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float32258 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float32
@@ -272,7 +273,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2 ALL
272add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2.o273add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2.o
273 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2.o274 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2.o
274 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0275 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
275- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec276+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
276 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false277 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
277 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__278 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
278 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_bfloat16279 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_bfloat16
@@ -365,6 +366,7 @@ add_custom_target(hccl_aiv_all_reduce_op_910_96 ALL
365 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_96.o)366 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_96.o)
366add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_96.o367add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_96.o
367 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec368 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_reduce_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
369+ -DASCENDC_DUMP=${ASCENDC_DUMP}
368 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000370 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
369 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/371 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
370 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw372 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
@@ -384,7 +386,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_int8_v2_910_96 ALL
384add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2_910_96.o386add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2_910_96.o
385 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2_910_96.o387 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int8_v2_910_96.o
386 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0388 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
387- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec389+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
388 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false390 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
389 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__391 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
390 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int8392 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int8
@@ -406,7 +408,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_int16_v2_910_96 ALL
406add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2_910_96.o408add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2_910_96.o
407 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2_910_96.o409 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int16_v2_910_96.o
408 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0410 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
409- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec411+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
410 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false412 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
411 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__413 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
412 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int16414 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int16
@@ -428,7 +430,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_int32_v2_910_96 ALL
428add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2_910_96.o430add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2_910_96.o
429 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2_910_96.o431 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_int32_v2_910_96.o
430 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0432 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
431- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec433+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
432 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false434 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
433 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__435 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
434 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int32436 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_int32
@@ -450,7 +452,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_float16_v2_910_96 ALL
450add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2_910_96.o452add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2_910_96.o
451 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2_910_96.o453 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float16_v2_910_96.o
452 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0454 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
453- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec455+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
454 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false456 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
455 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__457 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
456 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float16458 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float16
@@ -472,7 +474,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_float32_v2_910_96 ALL
472add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2_910_96.o474add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2_910_96.o
473 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2_910_96.o475 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_float32_v2_910_96.o
474 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0476 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
475- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec477+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
476 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false478 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
477 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__479 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
478 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float32480 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_float32
@@ -494,7 +496,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2_910_96 ALL
494add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2_910_96.o496add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2_910_96.o
495 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2_910_96.o497 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_oneshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_oneshot_bfloat16_v2_910_96.o
496 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0498 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
497- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec499+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
498 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false500 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
499 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__501 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
500 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_bfloat16502 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_oneshot=sk_allreduce_mesh_1d_oneshot_bfloat16
@@ -516,7 +518,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_int8_v2_910_96 ALL
516add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2_910_96.o518add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2_910_96.o
517 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2_910_96.o519 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int8_v2_910_96.o
518 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0520 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
519- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec521+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
520 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false522 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
521 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__523 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
522 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int8524 -DHCCL_DTYPE_INT8 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int8
@@ -538,7 +540,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_int16_v2_910_96 ALL
538add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2_910_96.o540add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2_910_96.o
539 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2_910_96.o541 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int16_v2_910_96.o
540 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0542 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
541- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec543+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
542 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false544 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
543 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__545 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
544 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int16546 -DHCCL_DTYPE_INT16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int16
@@ -560,7 +562,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_int32_v2_910_96 ALL
560add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2_910_96.o562add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2_910_96.o
561 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2_910_96.o563 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_int32_v2_910_96.o
562 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0564 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
563- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec565+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
564 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false566 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
565 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__567 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
566 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int32568 -DHCCL_DTYPE_INT32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_int32
@@ -582,7 +584,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_float16_v2_910_96 ALL
582add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2_910_96.o584add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2_910_96.o
583 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2_910_96.o585 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float16_v2_910_96.o
584 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0586 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
585- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec587+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
586 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false588 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
587 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__589 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
588 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float16590 -DHCCL_DTYPE_FP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float16
@@ -604,7 +606,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_float32_v2_910_96 ALL
604add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2_910_96.o606add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2_910_96.o
605 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2_910_96.o607 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_float32_v2_910_96.o
606 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0608 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
607- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec609+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
608 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false610 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
609 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__611 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
610 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float32612 -DHCCL_DTYPE_FP32 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_float32
@@ -626,7 +628,7 @@ add_custom_target(hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2_910_96 ALL
626add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2_910_96.o628add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2_910_96.o
627 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2_910_96.o629 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_ar_mesh_1d_twoshot_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_ar_superkernel_mesh_1d_twoshot_bfloat16_v2_910_96.o
628 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0630 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
629- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec631+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
630 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false632 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
631 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__633 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
632 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_bfloat16634 -DHCCL_DTYPE_BFP16 -Dsk_allreduce_mesh_1d_twoshot=sk_allreduce_mesh_1d_twoshot_bfloat16
Msrc/ops/all_reduce/template/aiv/kernel/aiv_all_reduce_op.h+2-0
@@ -20,6 +20,7 @@ using namespace AscendC;
20#define AIV_ALL_REDUCE_ONESHOT_KERNEL_BATCH_DEF(type) \20#define AIV_ALL_REDUCE_ONESHOT_KERNEL_BATCH_DEF(type) \
21 extern "C" __global__ __aicore__ void aiv_allreduce_##type(KERNEL_ARGS_DEF) \21 extern "C" __global__ __aicore__ void aiv_allreduce_##type(KERNEL_ARGS_DEF) \
22 { \22 { \
23+ AIV_INFO_HINT; \
23 return AivAllReduceV2Mesh1DOneShot<type>(KERNEL_ARGS_CALL); \24 return AivAllReduceV2Mesh1DOneShot<type>(KERNEL_ARGS_CALL); \
24 } \25 } \
25 EXPORT_AIV_META_INFO(aiv_allreduce_##type)26 EXPORT_AIV_META_INFO(aiv_allreduce_##type)
@@ -27,6 +28,7 @@ using namespace AscendC;
27#define AIV_ALLREDUCE_MESH1D_TWOSHOT_KERNEL_BATCH_DEF(type) \28#define AIV_ALLREDUCE_MESH1D_TWOSHOT_KERNEL_BATCH_DEF(type) \
28 extern "C" __global__ __aicore__ void aiv_allreduce_mesh1d_twoshot_##type(KERNEL_ARGS_DEF) \29 extern "C" __global__ __aicore__ void aiv_allreduce_mesh1d_twoshot_##type(KERNEL_ARGS_DEF) \
29 { \30 { \
31+ AIV_INFO_HINT; \
30 return AivAllReduceV2Mesh1DTwoShot<type>(KERNEL_ARGS_CALL); \32 return AivAllReduceV2Mesh1DTwoShot<type>(KERNEL_ARGS_CALL); \
31 } \33 } \
32 EXPORT_AIV_META_INFO(aiv_allreduce_mesh1d_twoshot_##type)34 EXPORT_AIV_META_INFO(aiv_allreduce_mesh1d_twoshot_##type)
Msrc/ops/all_reduce/template/aiv/kernel/aiv_ar_mesh_1d_oneshot_superkernel.h+1-0
@@ -13,6 +13,7 @@
13 13 
14extern "C" __aicore__ void sk_allreduce_mesh_1d_oneshot(SUPERKERNEL_LITE_ARGS_DEF)14extern "C" __aicore__ void sk_allreduce_mesh_1d_oneshot(SUPERKERNEL_LITE_ARGS_DEF)
15{15{
16+ AIV_INFO_HINT;
16 SUPERKERNEL_LITE_ARGS_EXTRACT;17 SUPERKERNEL_LITE_ARGS_EXTRACT;
17 return sk_ar_mesh_1d_oneshot(SUPERKERNEL_ARGS_CALL);18 return sk_ar_mesh_1d_oneshot(SUPERKERNEL_ARGS_CALL);
18}19}
Msrc/ops/all_reduce/template/aiv/kernel/aiv_ar_mesh_1d_twoshot_superkernel.h+1-0
@@ -13,6 +13,7 @@
13 13 
14extern "C" __aicore__ void sk_allreduce_mesh_1d_twoshot(SUPERKERNEL_LITE_ARGS_DEF)14extern "C" __aicore__ void sk_allreduce_mesh_1d_twoshot(SUPERKERNEL_LITE_ARGS_DEF)
15{15{
16+ AIV_INFO_HINT;
16 SUPERKERNEL_LITE_ARGS_EXTRACT;17 SUPERKERNEL_LITE_ARGS_EXTRACT;
17 return sk_ar_mesh_1d_twoshot(SUPERKERNEL_ARGS_CALL);18 return sk_ar_mesh_1d_twoshot(SUPERKERNEL_ARGS_CALL);
18}19}
Msrc/ops/all_to_all_v/template/aiv/kernel/CMakeLists.txt+26-22
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_all_to_all_v_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_v_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_v_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -34,6 +35,7 @@ add_custom_target(hccl_aiv_all_to_all_op_910_95 ALL
34 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_95.o)35 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_95.o)
35add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_95.o36add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_95.o
36 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec37 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
38+ -DASCENDC_DUMP=${ASCENDC_DUMP}
37 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800039 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
38 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"40 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
39 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"41 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -64,7 +66,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int8_v2 ALL
64add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2.o66add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2.o
65 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2.o67 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2.o
66 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=068 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
67- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec69+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
68 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false70 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
69 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__71 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
70 -DHCCL_DTYPE_INT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int872 -DHCCL_DTYPE_INT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int8
@@ -85,7 +87,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int16_v2 ALL
85add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2.o87add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2.o
86 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2.o88 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2.o
87 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=089 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
88- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec90+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
89 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false91 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
90 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__92 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
91 -DHCCL_DTYPE_INT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int1693 -DHCCL_DTYPE_INT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int16
@@ -106,7 +108,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int32_v2 ALL
106add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2.o108add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2.o
107 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2.o109 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2.o
108 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0110 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
109- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec111+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
110 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false112 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
111 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__113 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
112 -DHCCL_DTYPE_INT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int32114 -DHCCL_DTYPE_INT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int32
@@ -127,7 +129,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_float16_v2 ALL
127add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2.o129add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2.o
128 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2.o130 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2.o
129 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0131 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
130- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec132+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
131 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false133 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
132 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__134 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
133 -DHCCL_DTYPE_FP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float16135 -DHCCL_DTYPE_FP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float16
@@ -148,7 +150,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_float32_v2 ALL
148add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2.o150add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2.o
149 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2.o151 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2.o
150 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0152 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
151- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec153+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
152 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false154 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
153 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__155 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
154 -DHCCL_DTYPE_FP32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float32156 -DHCCL_DTYPE_FP32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float32
@@ -169,7 +171,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_bfloat16_v2 ALL
169add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2.o171add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2.o
170 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2.o172 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2.o
171 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0173 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
172- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec174+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
173 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false175 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
174 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__176 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
175 -DHCCL_DTYPE_BFP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_bfloat16177 -DHCCL_DTYPE_BFP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_bfloat16
@@ -190,7 +192,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint8_v2 ALL
190add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2.o192add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2.o
191 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2.o193 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2.o
192 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0194 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
193- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec195+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
194 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false196 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
195 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__197 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
196 -DHCCL_DTYPE_UINT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint8198 -DHCCL_DTYPE_UINT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint8
@@ -211,7 +213,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint16_v2 ALL
211add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2.o213add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2.o
212 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2.o214 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2.o
213 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0215 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
214- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec216+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
215 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false217 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
216 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__218 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
217 -DHCCL_DTYPE_UINT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint16219 -DHCCL_DTYPE_UINT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint16
@@ -232,7 +234,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint32_v2 ALL
232add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2.o234add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2.o
233 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2.o235 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2.o
234 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0236 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
235- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec237+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
236 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false238 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
237 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__239 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
238 -DHCCL_DTYPE_UINT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint32240 -DHCCL_DTYPE_UINT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint32
@@ -253,7 +255,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int64_v2 ALL
253add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2.o255add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2.o
254 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2.o256 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2.o
255 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0257 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
256- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec258+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
257 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false259 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
258 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__260 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
259 -DHCCL_DTYPE_INT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int64261 -DHCCL_DTYPE_INT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int64
@@ -274,7 +276,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint64_v2 ALL
274add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2.o276add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2.o
275 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2.o277 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2.o
276 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0278 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
277- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec279+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
278 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false280 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
279 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__281 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
280 -DHCCL_DTYPE_UINT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint64282 -DHCCL_DTYPE_UINT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint64
@@ -361,6 +363,7 @@ add_custom_target(hccl_aiv_all_to_all_v_op_910_96 ALL
361 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_96.o)363 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_96.o)
362add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_96.o364add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_96.o
363 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_v_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec365 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_v_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_v_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
366+ -DASCENDC_DUMP=${ASCENDC_DUMP}
364 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000367 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
365 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/368 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
366 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw369 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
@@ -379,6 +382,7 @@ add_custom_target(hccl_aiv_all_to_all_op_910_96 ALL
379 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_96.o)382 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_96.o)
380add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_96.o383add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_96.o
381 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec384 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_all_to_all_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_all_to_all_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
385+ -DASCENDC_DUMP=${ASCENDC_DUMP}
382 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000386 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
383 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/387 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
384 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw388 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
@@ -410,7 +414,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int8_v2_910_96 ALL
410add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2_910_96.o414add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2_910_96.o
411 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2_910_96.o415 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int8_v2_910_96.o
412 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0416 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
413- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec417+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
414 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false418 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
415 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__419 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
416 -DHCCL_DTYPE_INT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int8420 -DHCCL_DTYPE_INT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int8
@@ -432,7 +436,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int16_v2_910_96 ALL
432add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2_910_96.o436add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2_910_96.o
433 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2_910_96.o437 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int16_v2_910_96.o
434 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0438 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
435- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec439+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
436 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false440 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
437 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__441 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
438 -DHCCL_DTYPE_INT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int16442 -DHCCL_DTYPE_INT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int16
@@ -454,7 +458,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int32_v2_910_96 ALL
454add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2_910_96.o458add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2_910_96.o
455 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2_910_96.o459 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int32_v2_910_96.o
456 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0460 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
457- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec461+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
458 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false462 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
459 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__463 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
460 -DHCCL_DTYPE_INT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int32464 -DHCCL_DTYPE_INT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int32
@@ -476,7 +480,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_float16_v2_910_96 ALL
476add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2_910_96.o480add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2_910_96.o
477 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2_910_96.o481 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float16_v2_910_96.o
478 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0482 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
479- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec483+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
480 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false484 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
481 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__485 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
482 -DHCCL_DTYPE_FP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float16486 -DHCCL_DTYPE_FP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float16
@@ -498,7 +502,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_float32_v2_910_96 ALL
498add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2_910_96.o502add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2_910_96.o
499 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2_910_96.o503 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_float32_v2_910_96.o
500 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0504 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
501- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec505+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
502 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false506 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
503 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__507 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
504 -DHCCL_DTYPE_FP32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float32508 -DHCCL_DTYPE_FP32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_float32
@@ -520,7 +524,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_bfloat16_v2_910_96 ALL
520add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2_910_96.o524add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2_910_96.o
521 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2_910_96.o525 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_bfloat16_v2_910_96.o
522 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0526 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
523- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec527+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
524 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false528 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
525 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__529 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
526 -DHCCL_DTYPE_BFP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_bfloat16530 -DHCCL_DTYPE_BFP16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_bfloat16
@@ -542,7 +546,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint8_v2_910_96 ALL
542add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2_910_96.o546add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2_910_96.o
543 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2_910_96.o547 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint8_v2_910_96.o
544 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0548 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
545- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec549+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
546 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false550 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
547 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__551 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
548 -DHCCL_DTYPE_UINT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint8552 -DHCCL_DTYPE_UINT8 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint8
@@ -564,7 +568,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint16_v2_910_96 ALL
564add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2_910_96.o568add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2_910_96.o
565 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2_910_96.o569 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint16_v2_910_96.o
566 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0570 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
567- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec571+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
568 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false572 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
569 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__573 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
570 -DHCCL_DTYPE_UINT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint16574 -DHCCL_DTYPE_UINT16 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint16
@@ -586,7 +590,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint32_v2_910_96 ALL
586add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2_910_96.o590add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2_910_96.o
587 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2_910_96.o591 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint32_v2_910_96.o
588 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0592 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
589- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec593+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
590 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false594 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
591 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__595 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
592 -DHCCL_DTYPE_UINT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint32596 -DHCCL_DTYPE_UINT32 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint32
@@ -608,7 +612,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_int64_v2_910_96 ALL
608add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2_910_96.o612add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2_910_96.o
609 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2_910_96.o613 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_int64_v2_910_96.o
610 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0614 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
611- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec615+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
612 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false616 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
613 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__617 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
614 -DHCCL_DTYPE_INT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int64618 -DHCCL_DTYPE_INT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_int64
@@ -630,7 +634,7 @@ add_custom_target(hccl_a2a_superkernel_mesh_1d_uint64_v2_910_96 ALL
630add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2_910_96.o634add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2_910_96.o
631 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2_910_96.o635 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_a2a_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_a2a_superkernel_mesh_1d_uint64_v2_910_96.o
632 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0636 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
633- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec637+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
634 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false638 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
635 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__639 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
636 -DHCCL_DTYPE_UINT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint64640 -DHCCL_DTYPE_UINT64 -Dsk_alltoall_mesh_1d=sk_alltoall_mesh_1d_uint64
Msrc/ops/all_to_all_v/template/aiv/kernel/aiv_a2a_mesh_1d_superkernel.h+1-0
@@ -13,6 +13,7 @@
13 13 
14extern "C" __aicore__ void sk_alltoall_mesh_1d(SUPERKERNEL_LITE_ARGS_DEF)14extern "C" __aicore__ void sk_alltoall_mesh_1d(SUPERKERNEL_LITE_ARGS_DEF)
15{15{
16+ AIV_INFO_HINT;
16 SUPERKERNEL_LITE_ARGS_EXTRACT;17 SUPERKERNEL_LITE_ARGS_EXTRACT;
17 return sk_a2a_mesh_1d(SUPERKERNEL_ARGS_CALL);18 return sk_a2a_mesh_1d(SUPERKERNEL_ARGS_CALL);
18}19}
Msrc/ops/all_to_all_v/template/aiv/kernel/aiv_all_to_all_op.h+1-0
@@ -18,6 +18,7 @@ using namespace AscendC;
18#define AIV_ALL_TO_ALL_KERNEL_BATCH_DEF(type) \18#define AIV_ALL_TO_ALL_KERNEL_BATCH_DEF(type) \
19 extern "C" __global__ __aicore__ void aiv_alltoall_##type(KERNEL_ARGS_DEF) \19 extern "C" __global__ __aicore__ void aiv_alltoall_##type(KERNEL_ARGS_DEF) \
20 { \20 { \
21+ AIV_INFO_HINT; \
21 AivAlltoAllV2Mesh1D<type>(KERNEL_ARGS_CALL); \22 AivAlltoAllV2Mesh1D<type>(KERNEL_ARGS_CALL); \
22 } \23 } \
23 EXPORT_AIV_META_INFO(aiv_alltoall_##type)24 EXPORT_AIV_META_INFO(aiv_alltoall_##type)
Msrc/ops/all_to_all_v/template/aiv/kernel/aiv_all_to_all_v_op.h+1-0
@@ -18,6 +18,7 @@ using namespace AscendC;
18#define AIV_ALL_TO_ALL_V_KERNEL_BATCH_DEF(type) \18#define AIV_ALL_TO_ALL_V_KERNEL_BATCH_DEF(type) \
19 extern "C" __global__ __aicore__ void aiv_alltoallv_##type(EXTERN_KERNEL_ARGS_DEF_V2) \19 extern "C" __global__ __aicore__ void aiv_alltoallv_##type(EXTERN_KERNEL_ARGS_DEF_V2) \
20 { \20 { \
21+ AIV_INFO_HINT; \
21 AivAlltoAllVV2Mesh1D<type>(EXTERN_KERNEL_ARGS_CALL); \22 AivAlltoAllVV2Mesh1D<type>(EXTERN_KERNEL_ARGS_CALL); \
22 } \23 } \
23 EXPORT_AIV_META_INFO(aiv_alltoallv_##type)24 EXPORT_AIV_META_INFO(aiv_alltoallv_##type)
Msrc/ops/broadcast/template/aiv/kernel/CMakeLists.txt+2-0
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_broadcast_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_broadcast_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_broadcast_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,6 +42,7 @@ add_custom_target(hccl_aiv_broadcast_op_910_96 ALL
41 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_96.o)42 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_96.o)
42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_96.o43add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_96.o
43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_broadcast_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec44 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_broadcast_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_broadcast_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
45+ -DASCENDC_DUMP=${ASCENDC_DUMP}
44 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800046 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
45 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/47 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
46 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw48 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
Msrc/ops/broadcast/template/aiv/kernel/aiv_broadcast_op.h+1-0
@@ -19,6 +19,7 @@ using namespace AscendC;
19#define AIV_BROADCAST_KERNEL_BATCH_DEF(type) \19#define AIV_BROADCAST_KERNEL_BATCH_DEF(type) \
20 extern "C" __global__ __aicore__ void aiv_broadcast_##type(KERNEL_ARGS_DEF) \20 extern "C" __global__ __aicore__ void aiv_broadcast_##type(KERNEL_ARGS_DEF) \
21 { \21 { \
22+ AIV_INFO_HINT; \
22 return AivBroadcastV2Mesh1D<type>(KERNEL_ARGS_CALL); \23 return AivBroadcastV2Mesh1D<type>(KERNEL_ARGS_CALL); \
23 } \24 } \
24 EXPORT_AIV_META_INFO(aiv_broadcast_##type)25 EXPORT_AIV_META_INFO(aiv_broadcast_##type)
Msrc/ops/op_common/template/aiv/aiv_defines.h+7-0
@@ -156,4 +156,11 @@ enum class CommPattern {
156 156 
157#define SUPERKERNEL_CLASS_INIT hiddenInput, input, output157#define SUPERKERNEL_CLASS_INIT hiddenInput, input, output
158 158 
159+#define AIV_INFO(format, ...) \
160+ do { \
161+ AscendC::PRINTF(format, ##__VA_ARGS__); \
162+ } while (0)
163+ 
164+#define AIV_INFO_HINT AIV_INFO("Aiv log dump is enabled in %s\n", __func__)
165+ 
159#endif // AIV_DEFINES_H166#endif // AIV_DEFINES_H
Msrc/ops/recv/template/aiv/kernel/CMakeLists.txt+2-0
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_recv_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_recv_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_recv_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,6 +42,7 @@ add_custom_target(hccl_aiv_recv_op_910_96 ALL
41 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_96.o)42 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_96.o)
42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_96.o43add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_96.o
43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_recv_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec44 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_recv_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_recv_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
45+ -DASCENDC_DUMP=${ASCENDC_DUMP}
44 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800046 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
45 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/47 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
46 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw48 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
Msrc/ops/recv/template/aiv/kernel/aiv_recv_op.h+1-0
@@ -18,6 +18,7 @@ using namespace AscendC;
18#define AIV_RECV_KERNEL_BATCH_DEF(type) \18#define AIV_RECV_KERNEL_BATCH_DEF(type) \
19 extern "C" __global__ __aicore__ void aiv_recv_##type(KERNEL_ARGS_DEF) \19 extern "C" __global__ __aicore__ void aiv_recv_##type(KERNEL_ARGS_DEF) \
20 { \20 { \
21+ AIV_INFO_HINT; \
21 AivRecvV2Mesh1D<type>(KERNEL_ARGS_CALL); \22 AivRecvV2Mesh1D<type>(KERNEL_ARGS_CALL); \
22 } \23 } \
23 EXPORT_AIV_META_INFO(aiv_recv_##type)24 EXPORT_AIV_META_INFO(aiv_recv_##type)
Msrc/ops/reduce/template/aiv/kernel/CMakeLists.txt+2-0
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_reduce_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,6 +42,7 @@ add_custom_target(hccl_aiv_reduce_op_910_96 ALL
41 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_96.o)42 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_96.o)
42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_96.o43add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_96.o
43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec44 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
45+ -DASCENDC_DUMP=${ASCENDC_DUMP}
44 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800046 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
45 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/47 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
46 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw48 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
Msrc/ops/reduce/template/aiv/kernel/aiv_reduce_op.h+1-0
@@ -19,6 +19,7 @@ using namespace AscendC;
19#define AIV_REDUCE_KERNEL_BATCH_DEF(type) \19#define AIV_REDUCE_KERNEL_BATCH_DEF(type) \
20 extern "C" __global__ __aicore__ void aiv_reduce_##type(KERNEL_ARGS_DEF) \20 extern "C" __global__ __aicore__ void aiv_reduce_##type(KERNEL_ARGS_DEF) \
21 { \21 { \
22+ AIV_INFO_HINT; \
22 return AivReduceV2Mesh1D<type>(KERNEL_ARGS_CALL); \23 return AivReduceV2Mesh1D<type>(KERNEL_ARGS_CALL); \
23 } \24 } \
24 EXPORT_AIV_META_INFO(aiv_reduce_##type)25 EXPORT_AIV_META_INFO(aiv_reduce_##type)
Msrc/ops/reduce_scatter/template/aiv/kernel/CMakeLists.txt+14-12
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_reduce_scatter_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,7 +42,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_int8_v2 ALL
41add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2.o42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2.o
42 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2.o43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2.o
43 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=044 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
44- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec45+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
45 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false46 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
46 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__47 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
47 -DHCCL_DTYPE_INT8 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int848 -DHCCL_DTYPE_INT8 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int8
@@ -62,7 +63,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_int16_v2 ALL
62add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2.o63add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2.o
63 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2.o64 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2.o
64 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=065 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
65- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec66+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
66 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false67 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
67 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__68 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
68 -DHCCL_DTYPE_INT16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int1669 -DHCCL_DTYPE_INT16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int16
@@ -83,7 +84,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_int32_v2 ALL
83add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2.o84add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2.o
84 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2.o85 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2.o
85 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=086 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
86- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec87+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
87 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false88 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
88 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__89 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
89 -DHCCL_DTYPE_INT32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int3290 -DHCCL_DTYPE_INT32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int32
@@ -104,7 +105,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_float16_v2 ALL
104add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2.o105add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2.o
105 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2.o106 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2.o
106 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0107 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
107- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec108+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
108 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false109 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
109 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__110 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
110 -DHCCL_DTYPE_FP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float16111 -DHCCL_DTYPE_FP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float16
@@ -125,7 +126,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_float32_v2 ALL
125add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2.o126add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2.o
126 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2.o127 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2.o
127 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0128 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
128- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec129+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
129 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false130 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
130 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__131 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
131 -DHCCL_DTYPE_FP32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float32132 -DHCCL_DTYPE_FP32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float32
@@ -146,7 +147,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_bfloat16_v2 ALL
146add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2.o147add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2.o
147 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2.o148 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2.o
148 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0149 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
149- -DASCENDC_DUMP --cce-aicore-arch=dav-c310-vec150+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-c310-vec
150 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false151 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
151 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__152 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=310 -D__DAV_C310__
152 -DHCCL_DTYPE_BFP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_bfloat16153 -DHCCL_DTYPE_BFP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_bfloat16
@@ -203,6 +204,7 @@ add_custom_target(hccl_aiv_reduce_scatter_op_910_96 ALL
203 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_96.o)204 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_96.o)
204add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_96.o205add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_96.o
205 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec206 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_reduce_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_reduce_scatter_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
207+ -DASCENDC_DUMP=${ASCENDC_DUMP}
206 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000208 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
207 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/209 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
208 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw210 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
@@ -228,7 +230,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_int8_v2_910_96 ALL
228add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2_910_96.o230add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2_910_96.o
229 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2_910_96.o231 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int8_v2_910_96.o
230 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0232 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
231- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec233+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
232 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false234 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
233 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__235 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
234 -DHCCL_DTYPE_INT8 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int8236 -DHCCL_DTYPE_INT8 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int8
@@ -250,7 +252,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_int16_v2_910_96 ALL
250add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2_910_96.o252add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2_910_96.o
251 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2_910_96.o253 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int16_v2_910_96.o
252 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0254 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
253- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec255+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
254 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false256 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
255 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__257 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
256 -DHCCL_DTYPE_INT16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int16258 -DHCCL_DTYPE_INT16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int16
@@ -272,7 +274,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_int32_v2_910_96 ALL
272add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2_910_96.o274add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2_910_96.o
273 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2_910_96.o275 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_int32_v2_910_96.o
274 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0276 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
275- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec277+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
276 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false278 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
277 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__279 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
278 -DHCCL_DTYPE_INT32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int32280 -DHCCL_DTYPE_INT32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_int32
@@ -294,7 +296,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_float16_v2_910_96 ALL
294add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2_910_96.o296add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2_910_96.o
295 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2_910_96.o297 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float16_v2_910_96.o
296 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0298 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
297- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec299+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
298 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false300 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
299 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__301 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
300 -DHCCL_DTYPE_FP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float16302 -DHCCL_DTYPE_FP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float16
@@ -316,7 +318,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_float32_v2_910_96 ALL
316add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2_910_96.o318add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2_910_96.o
317 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2_910_96.o319 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_float32_v2_910_96.o
318 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0320 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
319- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec321+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
320 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false322 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
321 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__323 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
322 -DHCCL_DTYPE_FP32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float32324 -DHCCL_DTYPE_FP32 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_float32
@@ -338,7 +340,7 @@ add_custom_target(hccl_rs_superkernel_mesh_1d_bfloat16_v2_910_96 ALL
338add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2_910_96.o340add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2_910_96.o
339 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2_910_96.o341 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_rs_mesh_1d_superkernel.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_rs_superkernel_mesh_1d_bfloat16_v2_910_96.o
340 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0342 -DTILING_KEY_VAR=0UL -D__ASCENDC_ENABLE_SUPER_KERNEL__ -DL2_CACHE_HINT -DCONST_TILING -DDETERMINISTIC_MODE=0
341- -DASCENDC_DUMP --cce-aicore-arch=dav-920r1-vec343+ -DASCENDC_DUMP=${ASCENDC_DUMP} --cce-aicore-arch=dav-920r1-vec
342 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false344 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-dcci-insert-for-scalar=false
343 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__345 -mllvm -cce-aicore-stack-size=0x8000 -D__CCE_AICORE__=920 -D__DAV_920R1__
344 -DHCCL_DTYPE_BFP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_bfloat16346 -DHCCL_DTYPE_BFP16 -Dsk_reducescatter_mesh_1d=sk_reducescatter_mesh_1d_bfloat16
Msrc/ops/reduce_scatter/template/aiv/kernel/aiv_reduce_scatter_op.h+1-0
@@ -23,6 +23,7 @@ using namespace AscendC;
23#define AIV_REDUCE_SCATTER_KERNEL_BATCH_DEF(type) \23#define AIV_REDUCE_SCATTER_KERNEL_BATCH_DEF(type) \
24 extern "C" __global__ __aicore__ void aiv_reduce_scatter_##type(KERNEL_ARGS_DEF) \24 extern "C" __global__ __aicore__ void aiv_reduce_scatter_##type(KERNEL_ARGS_DEF) \
25 { \25 { \
26+ AIV_INFO_HINT; \
26 if (numBlocks > 2 * rankSize) { \27 if (numBlocks > 2 * rankSize) { \
27 AivReduceScatterV2Mesh1DBigData<type>(KERNEL_ARGS_CALL); \28 AivReduceScatterV2Mesh1DBigData<type>(KERNEL_ARGS_CALL); \
28 } else if (numBlocks >= rankSize) { \29 } else if (numBlocks >= rankSize) { \
Msrc/ops/reduce_scatter/template/aiv/kernel/aiv_rs_mesh_1d_superkernel.h+1-0
@@ -48,6 +48,7 @@ __aicore__ inline void sk_rs_superkernel_dispatch(SUPERKERNEL_ARGS_DEF)
48 48 
49extern "C" __aicore__ void sk_reducescatter_mesh_1d(SUPERKERNEL_LITE_ARGS_DEF)49extern "C" __aicore__ void sk_reducescatter_mesh_1d(SUPERKERNEL_LITE_ARGS_DEF)
50{50{
51+ AIV_INFO_HINT;
51 SUPERKERNEL_LITE_ARGS_EXTRACT;52 SUPERKERNEL_LITE_ARGS_EXTRACT;
52 return sk_rs_superkernel_dispatch(SUPERKERNEL_ARGS_CALL);53 return sk_rs_superkernel_dispatch(SUPERKERNEL_ARGS_CALL);
53}54}
Msrc/ops/scatter/template/aiv/kernel/CMakeLists.txt+2-0
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_scatter_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,6 +42,7 @@ add_custom_target(hccl_aiv_scatter_op_910_96 ALL
41 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_96.o)42 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_96.o)
42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_96.o43add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_96.o
43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec44 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_scatter_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_scatter_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
45+ -DASCENDC_DUMP=${ASCENDC_DUMP}
44 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800046 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
45 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/47 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
46 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw48 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
Msrc/ops/scatter/template/aiv/kernel/aiv_scatter_op.h+1-0
@@ -19,6 +19,7 @@ using namespace AscendC;
19#define AIV_SCATTER_KERNEL_BATCH_DEF(type) \19#define AIV_SCATTER_KERNEL_BATCH_DEF(type) \
20 extern "C" __global__ __aicore__ void aiv_scatter_##type(KERNEL_ARGS_DEF) \20 extern "C" __global__ __aicore__ void aiv_scatter_##type(KERNEL_ARGS_DEF) \
21 { \21 { \
22+ AIV_INFO_HINT; \
22 return AivScatterV2Mesh1D<type>(KERNEL_ARGS_CALL); \23 return AivScatterV2Mesh1D<type>(KERNEL_ARGS_CALL); \
23 } \24 } \
24 EXPORT_AIV_META_INFO(aiv_scatter_##type)25 EXPORT_AIV_META_INFO(aiv_scatter_##type)
Msrc/ops/send/template/aiv/kernel/CMakeLists.txt+2-0
@@ -17,6 +17,7 @@ add_custom_target(hccl_aiv_send_op_910_95 ALL
17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_95.o)17 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_95.o)
18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_95.o18add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_95.o
19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_send_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec19 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_send_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_95_tmp.o --cce-aicore-arch=dav-c310-vec
20+ -DASCENDC_DUMP=${ASCENDC_DUMP}
20 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800021 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
21 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"22 "-I$<JOIN:$<TARGET_PROPERTY:asc_kernel_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
22 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"23 "-I$<JOIN:$<TARGET_PROPERTY:kernel_tiling_headers,INTERFACE_INCLUDE_DIRECTORIES>,;-I>"
@@ -41,6 +42,7 @@ add_custom_target(hccl_aiv_send_op_910_96 ALL
41 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_96.o)42 DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_96.o)
42add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_96.o43add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_96.o
43 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_send_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec44 COMMAND ${CCE_DIR}/ccec -c -x cce -O2 ${CMAKE_CURRENT_SOURCE_DIR}/aiv_send_op.h -o ${CMAKE_CURRENT_BINARY_DIR}/hccl_aiv_send_op_910_96_tmp.o --cce-aicore-arch=dav-920r1-vec
45+ -DASCENDC_DUMP=${ASCENDC_DUMP}
44 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x800046 --cce-aicore-input-parameter-size=28000 --cce-aicore-only -mllvm -cce-aicore-function-stack-size=0x8000 -mllvm -cce-aicore-stack-size=0x8000
45 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/47 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw/lib/
46 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw48 -I${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/tikcfw
Msrc/ops/send/template/aiv/kernel/aiv_send_op.h+1-0
@@ -18,6 +18,7 @@ using namespace AscendC;
18#define AIV_SEND_KERNEL_BATCH_DEF(type) \18#define AIV_SEND_KERNEL_BATCH_DEF(type) \
19 extern "C" __global__ __aicore__ void aiv_send_##type(KERNEL_ARGS_DEF) \19 extern "C" __global__ __aicore__ void aiv_send_##type(KERNEL_ARGS_DEF) \
20 { \20 { \
21+ AIV_INFO_HINT; \
21 AivSendV2Mesh1D<type>(KERNEL_ARGS_CALL); \22 AivSendV2Mesh1D<type>(KERNEL_ARGS_CALL); \
22 } \23 } \
23 EXPORT_AIV_META_INFO(aiv_send_##type)24 EXPORT_AIV_META_INFO(aiv_send_##type)