add <<<>>> kernel launch-asc-comm-AtomGit

Mexamples/ccu/ccu_all_to_all/op_host/alg_resource.cc+15-15

	@@ -22,33 +22,33 @@ namespace ops_hccl_a2a {
22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
23		23
24	static HcclResult AllocThreadAndChannelResource(	24	static HcclResult AllocThreadAndChannelResource(
25	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	25	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
26	uint32_t& kernelDie)	26	uint32_t& kernelDie)
27	{	27	{
28	ThreadHandle thread;	28	ThreadHandle thread;
29	constexpr uint32_t notifyNumOnMainThread = 0;	29	constexpr uint32_t notifyNumOnMainThread = 0;
30	RETURN_IF_HCCL_FAIL(	30	RETURN_IF_HCCL_FAIL(
31	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	31	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
32	resCtxHost.threads.push_back(thread);	32	resCtxHost.threads.push_back(thread);
33		33
34	- if (param_.rankSize == 1) {	34	+ if (opParam.rankSize == 1) {
35	return HCCL_SUCCESS;	35	return HCCL_SUCCESS;
36	}	36	}
37		37
38	- uint32_t channel_num = param_.rankSize - 1;	38	+ uint32_t channel_num = opParam.rankSize - 1;
39	kernelChannels.resize(channel_num);	39	kernelChannels.resize(channel_num);
40		40
41	uint32_t channelIndex = 0;	41	uint32_t channelIndex = 0;
42	bool hasKernelDie = false;	42	bool hasKernelDie = false;
43	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	43	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
44	- if (remoteRank == param_.myRank) {	44	+ if (remoteRank == opParam.myRank) {
45	continue;	45	continue;
46	}	46	}
47		47
48	uint32_t netLayer = 0, listSize = 0;	48	uint32_t netLayer = 0, listSize = 0;
49	CommLink* linkList = nullptr;	49	CommLink* linkList = nullptr;
50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(	50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(
51	- comm, netLayer, param_.myRank, remoteRank, &linkList,	51	+ comm, netLayer, opParam.myRank, remoteRank, &linkList,
52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息	52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息
53		53
54	HcclChannelDesc desc;	54	HcclChannelDesc desc;
	@@ -69,7 +69,7 @@ static HcclResult AllocThreadAndChannelResource(
69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
70	uint32_t channelDie = 0;	70	uint32_t channelDie = 0;
71	RETURN_IF_HCCL_FAIL(	71	RETURN_IF_HCCL_FAIL(
72	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	72	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
74	protocolExists = true;	74	protocolExists = true;
75	break;	75	break;
	@@ -88,7 +88,7 @@ static HcclResult AllocThreadAndChannelResource(
88	}	88	}
89		89
90	static HcclResult RegisterAllToAllKernel(	90	static HcclResult RegisterAllToAllKernel(
91	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	91	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
92	uint32_t kernelDie)	92	uint32_t kernelDie)
93	{	93	{
94	ccu_kernel_info kernelInfo;	94	ccu_kernel_info kernelInfo;
	@@ -99,8 +99,8 @@ static HcclResult RegisterAllToAllKernel(
99	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllToAllMesh1DMem2MemKernel);	99	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllToAllMesh1DMem2MemKernel);
100		100
101	auto kernel_arg = std::make_shared<CcuKernelArgAllToAllMesh1DMem2Mem>();	101	auto kernel_arg = std::make_shared<CcuKernelArgAllToAllMesh1DMem2Mem>();
102	- kernel_arg->rankSize = param_.rankSize;	102	+ kernel_arg->rankSize = opParam.rankSize;
103	- kernel_arg->rankId = param_.myRank;	103	+ kernel_arg->rankId = opParam.myRank;
104	kernelInfo.setKernelArg(kernel_arg);	104	kernelInfo.setKernelArg(kernel_arg);
105		105
106	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	106	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -139,17 +139,17 @@ static HcclResult RegisterAllToAllKernel(
139	return HCCL_SUCCESS;	139	return HCCL_SUCCESS;
140	}	140	}
141		141
142	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	142	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
143	{	143	{
144	std::vector<ChannelHandle> kernelChannels;	144	std::vector<ChannelHandle> kernelChannels;
145	uint32_t kernelDie = 0;	145	uint32_t kernelDie = 0;
146	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	146	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
147		147
148	- if (param_.rankSize == 1) {	148	+ if (opParam.rankSize == 1) {
149	return HCCL_SUCCESS;	149	return HCCL_SUCCESS;
150	}	150	}
151		151
152	- RETURN_IF_HCCL_FAIL(RegisterAllToAllKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	152	+ RETURN_IF_HCCL_FAIL(RegisterAllToAllKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
153	return HCCL_SUCCESS;	153	return HCCL_SUCCESS;
154	}	154	}
155		155

Mexamples/ccu/ccu_all_to_all/op_host/alg_resource.h+1-1

	@@ -37,7 +37,7 @@ struct OpParam {
37	uint32_t rankSize = 0;	37	uint32_t rankSize = 0;
38	};	38	};
39		39
40	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	40	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
41		41
42	} // namespace ops_hccl_a2a	42	} // namespace ops_hccl_a2a
43		43

Mexamples/ccu/ccu_all_to_all/op_host/exec_op.cc+12-17

	@@ -12,7 +12,6 @@
12	#include <algorithm>	12	#include <algorithm>
13	#include <ccu/hcomm/ccu_api_types.h>	13	#include <ccu/hcomm/ccu_api_types.h>
14	#include <hccl/hcomm_primitives.h>	14	#include <hccl/hcomm_primitives.h>
15	-#include <hcomm/hcomm_ccu_launch_api.h>
16	#include "ccu/hcomm/ccu_launch.h"	15	#include "ccu/hcomm/ccu_launch.h"
17	#include "ccu/hcomm/ccu_resource_api.h"	16	#include "ccu/hcomm/ccu_resource_api.h"
18	#include "alg_resource.h"	17	#include "alg_resource.h"
	@@ -83,8 +82,8 @@ static std::vector<uint64_t> CalGoSize(uint64_t size, const loop_group_config& c
83	}	82	}
84		83
85	static HcclResult LaunchCcuKernelSlice(	84	static HcclResult LaunchCcuKernelSlice(
86	- const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t inputToken, uint64_t outputToken,	85	+ const OpParam& opParam, const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t inputToken,
87	- uint64_t rankSliceSize, uint64_t sliceCount, uint64_t dataTypeSize)	86	+ uint64_t outputToken, uint64_t rankSliceSize, uint64_t sliceCount, uint64_t dataTypeSize)
88	{	87	{
89	uint64_t sliceSize = sliceCount * dataTypeSize;	88	uint64_t sliceSize = sliceCount * dataTypeSize;
90		89
	@@ -99,12 +98,8 @@ static HcclResult LaunchCcuKernelSlice(
99	sliceSize, goSize[0], goSize[1], goSize[2], goSize[3],	98	sliceSize, goSize[0], goSize[1], goSize[2], goSize[3],
100	};	99	};
101		100
102	- HcommCcuLaunchContextPod launch_context{};
103	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
104	- return HCCL_E_INTERNAL;
105	- }
106	CcuResult launchRet =	101	CcuResult launchRet =
107	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	102	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
108	if (launchRet != CCU_SUCCESS) {	103	if (launchRet != CCU_SUCCESS) {
109	return HCCL_E_INTERNAL;	104	return HCCL_E_INTERNAL;
110	}	105	}
	@@ -112,21 +107,21 @@ static HcclResult LaunchCcuKernelSlice(
112	return HCCL_SUCCESS;	107	return HCCL_SUCCESS;
113	}	108	}
114		109
115	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	110	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
116	{	111	{
117	constexpr uint64_t dataTypeSize = sizeof(float);	112	constexpr uint64_t dataTypeSize = sizeof(float);
118	- const uint64_t rankSliceSize = param_.perCount * dataTypeSize;	113	+ const uint64_t rankSliceSize = opParam.perCount * dataTypeSize;
119	- const uint64_t inputSize = rankSliceSize * param_.rankSize;	114	+ const uint64_t inputSize = rankSliceSize * opParam.rankSize;
120	const uint64_t outputSize = inputSize;	115	const uint64_t outputSize = inputSize;
121	- const uint64_t count_ = param_.perCount;	116	+ const uint64_t count_ = opParam.perCount;
122		117
123	if (count_ == 0) { // 数据量为 0，直接返回	118	if (count_ == 0) { // 数据量为 0，直接返回
124	return HcclResult::HCCL_SUCCESS;	119	return HcclResult::HCCL_SUCCESS;
125	}	120	}
126		121
127	- if (param_.rankSize == 1) { // 单卡场景直接本地拷贝	122	+ if (opParam.rankSize == 1) { // 单卡场景直接本地拷贝
128	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(	123	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(
129	- HcommLocalCopyOnThread(resCtx.threads[0], param_.outputPtr, param_.inputPtr, outputSize)));	124	+ HcommLocalCopyOnThread(resCtx.threads[0], opParam.outputPtr, opParam.inputPtr, outputSize)));
130	return HCCL_SUCCESS;	125	return HCCL_SUCCESS;
131	}	126	}
132		127
	@@ -138,8 +133,8 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
138		133
139	uint64_t inputToken = 0;	134	uint64_t inputToken = 0;
140	uint64_t outputToken = 0;	135	uint64_t outputToken = 0;
141	- uint64_t baseInputAddr = reinterpret_cast<uint64_t>(param_.inputPtr);	136	+ uint64_t baseInputAddr = reinterpret_cast<uint64_t>(opParam.inputPtr);
142	- uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(param_.outputPtr);	137	+ uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(opParam.outputPtr);
143	CcuResult tokenRet = asccomm_ccu_get_mem_token(baseInputAddr, inputSize, &inputToken);	138	CcuResult tokenRet = asccomm_ccu_get_mem_token(baseInputAddr, inputSize, &inputToken);
144	if (tokenRet != CCU_SUCCESS) {	139	if (tokenRet != CCU_SUCCESS) {
145	return HCCL_E_INTERNAL;	140	return HCCL_E_INTERNAL;
	@@ -155,7 +150,7 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
155	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;	150	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;
156		151
157	RETURN_IF_HCCL_FAIL(LaunchCcuKernelSlice(	152	RETURN_IF_HCCL_FAIL(LaunchCcuKernelSlice(
158	- resCtx, inputAddr, outputAddr, inputToken, outputToken, rankSliceSize, sliceCount, dataTypeSize));	153	+ opParam, resCtx, inputAddr, outputAddr, inputToken, outputToken, rankSliceSize, sliceCount, dataTypeSize));
159		154
160	processedDataCount += sliceCount;	155	processedDataCount += sliceCount;
161	}	156	}

Mexamples/ccu/ccu_all_to_all/op_host/exec_op.h+1-1

	@@ -15,7 +15,7 @@
15		15
16	namespace ops_hccl_a2a {	16	namespace ops_hccl_a2a {
17		17
18	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	18	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
19		19
20	} // namespace ops_hccl_a2a	20	} // namespace ops_hccl_a2a
21		21

Mexamples/ccu/ccu_all_to_all_vc/op_host/alg_resource.cc+15-15

	@@ -22,33 +22,33 @@ namespace ops_hccl_a2avc {
22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
23		23
24	static HcclResult AllocThreadAndChannelResource(	24	static HcclResult AllocThreadAndChannelResource(
25	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	25	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
26	uint32_t& kernelDie)	26	uint32_t& kernelDie)
27	{	27	{
28	ThreadHandle thread;	28	ThreadHandle thread;
29	constexpr uint32_t notifyNumOnMainThread = 0;	29	constexpr uint32_t notifyNumOnMainThread = 0;
30	RETURN_IF_HCCL_FAIL(	30	RETURN_IF_HCCL_FAIL(
31	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	31	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
32	resCtxHost.threads.push_back(thread);	32	resCtxHost.threads.push_back(thread);
33		33
34	- if (param_.rankSize == 1) {	34	+ if (opParam.rankSize == 1) {
35	return HCCL_SUCCESS;	35	return HCCL_SUCCESS;
36	}	36	}
37		37
38	- uint32_t channel_num = param_.rankSize - 1;	38	+ uint32_t channel_num = opParam.rankSize - 1;
39	kernelChannels.resize(channel_num);	39	kernelChannels.resize(channel_num);
40		40
41	uint32_t channelIndex = 0;	41	uint32_t channelIndex = 0;
42	bool hasKernelDie = false;	42	bool hasKernelDie = false;
43	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	43	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
44	- if (remoteRank == param_.myRank) {	44	+ if (remoteRank == opParam.myRank) {
45	continue;	45	continue;
46	}	46	}
47		47
48	uint32_t netLayer = 0, listSize = 0;	48	uint32_t netLayer = 0, listSize = 0;
49	CommLink* linkList = nullptr;	49	CommLink* linkList = nullptr;
50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(	50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(
51	- comm, netLayer, param_.myRank, remoteRank, &linkList,	51	+ comm, netLayer, opParam.myRank, remoteRank, &linkList,
52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息	52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息
53		53
54	HcclChannelDesc desc;	54	HcclChannelDesc desc;
	@@ -69,7 +69,7 @@ static HcclResult AllocThreadAndChannelResource(
69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
70	uint32_t channelDie = 0;	70	uint32_t channelDie = 0;
71	RETURN_IF_HCCL_FAIL(	71	RETURN_IF_HCCL_FAIL(
72	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	72	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
74	protocolExists = true;	74	protocolExists = true;
75	break;	75	break;
	@@ -88,7 +88,7 @@ static HcclResult AllocThreadAndChannelResource(
88	}	88	}
89		89
90	static HcclResult RegisterAllToAllVCKernel(	90	static HcclResult RegisterAllToAllVCKernel(
91	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	91	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
92	uint32_t kernelDie)	92	uint32_t kernelDie)
93	{	93	{
94	ccu_kernel_info kernelInfo;	94	ccu_kernel_info kernelInfo;
	@@ -99,8 +99,8 @@ static HcclResult RegisterAllToAllVCKernel(
99	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllToAllVCMesh1DKernel);	99	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllToAllVCMesh1DKernel);
100		100
101	auto kernel_arg = std::make_shared<CcuKernelArgAllToAllVCMesh1D>();	101	auto kernel_arg = std::make_shared<CcuKernelArgAllToAllVCMesh1D>();
102	- kernel_arg->rankSize = param_.rankSize;	102	+ kernel_arg->rankSize = opParam.rankSize;
103	- kernel_arg->rankId = param_.myRank;	103	+ kernel_arg->rankId = opParam.myRank;
104	kernelInfo.setKernelArg(kernel_arg);	104	kernelInfo.setKernelArg(kernel_arg);
105		105
106	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	106	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -139,17 +139,17 @@ static HcclResult RegisterAllToAllVCKernel(
139	return HCCL_SUCCESS;	139	return HCCL_SUCCESS;
140	}	140	}
141		141
142	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	142	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
143	{	143	{
144	std::vector<ChannelHandle> kernelChannels;	144	std::vector<ChannelHandle> kernelChannels;
145	uint32_t kernelDie = 0;	145	uint32_t kernelDie = 0;
146	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	146	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
147		147
148	- if (param_.rankSize == 1) {	148	+ if (opParam.rankSize == 1) {
149	return HCCL_SUCCESS;	149	return HCCL_SUCCESS;
150	}	150	}
151		151
152	- RETURN_IF_HCCL_FAIL(RegisterAllToAllVCKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	152	+ RETURN_IF_HCCL_FAIL(RegisterAllToAllVCKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
153	return HCCL_SUCCESS;	153	return HCCL_SUCCESS;
154	}	154	}
155		155

Mexamples/ccu/ccu_all_to_all_vc/op_host/alg_resource.h+1-1

	@@ -39,7 +39,7 @@ struct OpParam {
39	uint32_t rankSize = 0;	39	uint32_t rankSize = 0;
40	};	40	};
41		41
42	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	42	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
43		43
44	} // namespace ops_hccl_a2avc	44	} // namespace ops_hccl_a2avc
45		45

Mexamples/ccu/ccu_all_to_all_vc/op_host/exec_op.cc+16-21

	@@ -12,7 +12,6 @@
12	#include <vector>	12	#include <vector>
13	#include <ccu/hcomm/ccu_api_types.h>	13	#include <ccu/hcomm/ccu_api_types.h>
14	#include <hccl/hcomm_primitives.h>	14	#include <hccl/hcomm_primitives.h>
15	-#include <hcomm/hcomm_ccu_launch_api.h>
16	#include "ccu/hcomm/ccu_launch.h"	15	#include "ccu/hcomm/ccu_launch.h"
17	#include "ccu/hcomm/ccu_resource_api.h"	16	#include "ccu/hcomm/ccu_resource_api.h"
18	#include "alg_resource.h"	17	#include "alg_resource.h"
	@@ -32,50 +31,46 @@ static uint64_t GetBufferElements(const std::vector<uint64_t>& counts, const std
32	return elements;	31	return elements;
33	}	32	}
34		33
35	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	34	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
36	{	35	{
37	constexpr uint64_t dataTypeSize = sizeof(float);	36	constexpr uint64_t dataTypeSize = sizeof(float);
38	- const uint64_t inputSize = GetBufferElements(param_.sendCounts, param_.sendDispls) * dataTypeSize;	37	+ const uint64_t inputSize = GetBufferElements(opParam.sendCounts, opParam.sendDispls) * dataTypeSize;
39	- const uint64_t outputSize = GetBufferElements(param_.recvCounts, param_.recvDispls) * dataTypeSize;	38	+ const uint64_t outputSize = GetBufferElements(opParam.recvCounts, opParam.recvDispls) * dataTypeSize;
40	- if (param_.rankSize == 1) {	39	+ if (opParam.rankSize == 1) {
41	- if (param_.sendCounts[0] != param_.recvCounts[0]) {	40	+ if (opParam.sendCounts[0] != opParam.recvCounts[0]) {
42	return HCCL_E_PARA;	41	return HCCL_E_PARA;
43	}	42	}
44	- if (param_.sendCounts[0] == 0) {	43	+ if (opParam.sendCounts[0] == 0) {
45	return HCCL_SUCCESS;	44	return HCCL_SUCCESS;
46	}	45	}
47	- auto* src_ = static_cast<char>(param_.inputPtr) + param_.sendDispls[0] dataTypeSize;	46	+ auto* src_ = static_cast<char>(opParam.inputPtr) + opParam.sendDispls[0] dataTypeSize;
48	- auto* dst_ = static_cast<char>(param_.outputPtr) + param_.recvDispls[0] dataTypeSize;	47	+ auto* dst_ = static_cast<char>(opParam.outputPtr) + opParam.recvDispls[0] dataTypeSize;
49	return static_cast<HcclResult>(	48	return static_cast<HcclResult>(
50	- HcommLocalCopyOnThread(resCtx.threads[0], dst_, src_, param_.sendCounts[0] * dataTypeSize));	49	+ HcommLocalCopyOnThread(resCtx.threads[0], dst_, src_, opParam.sendCounts[0] * dataTypeSize));
51	}	50	}
52		51
53	uint64_t inputToken = 0;	52	uint64_t inputToken = 0;
54	uint64_t outputToken = 0;	53	uint64_t outputToken = 0;
55	- const uint64_t inputAddr = reinterpret_cast<uint64_t>(param_.inputPtr);	54	+ const uint64_t inputAddr = reinterpret_cast<uint64_t>(opParam.inputPtr);
56	- const uint64_t outputAddr = reinterpret_cast<uint64_t>(param_.outputPtr);	55	+ const uint64_t outputAddr = reinterpret_cast<uint64_t>(opParam.outputPtr);
57	if ((inputSize != 0 && asccomm_ccu_get_mem_token(inputAddr, inputSize, &inputToken) != CCU_SUCCESS) \|\|	56	if ((inputSize != 0 && asccomm_ccu_get_mem_token(inputAddr, inputSize, &inputToken) != CCU_SUCCESS) \|\|
58	(outputSize != 0 && asccomm_ccu_get_mem_token(outputAddr, outputSize, &outputToken) != CCU_SUCCESS)) {	57	(outputSize != 0 && asccomm_ccu_get_mem_token(outputAddr, outputSize, &outputToken) != CCU_SUCCESS)) {
59	return HCCL_E_INTERNAL;	58	return HCCL_E_INTERNAL;
60	}	59	}
61		60
62	std::vector<uint64_t> task_args = {inputAddr, outputAddr, inputToken, outputToken};	61	std::vector<uint64_t> task_args = {inputAddr, outputAddr, inputToken, outputToken};
63	- for (uint32_t rankIdx = 0; rankIdx < param_.rankSize; rankIdx++) {	62	+ for (uint32_t rankIdx = 0; rankIdx < opParam.rankSize; rankIdx++) {
64	- const uint64_t sendBytes = param_.sendCounts[rankIdx] * dataTypeSize;	63	+ const uint64_t sendBytes = opParam.sendCounts[rankIdx] * dataTypeSize;
65	if (sendBytes > MAX_TRANSFER_BYTES) {	64	if (sendBytes > MAX_TRANSFER_BYTES) {
66	return HCCL_E_NOT_SUPPORT;	65	return HCCL_E_NOT_SUPPORT;
67	}	66	}
68	task_args.push_back(sendBytes);	67	task_args.push_back(sendBytes);
69	- task_args.push_back(param_.sendDispls[rankIdx] * dataTypeSize);	68	+ task_args.push_back(opParam.sendDispls[rankIdx] * dataTypeSize);
70	- task_args.push_back(param_.recvDispls[rankIdx] * dataTypeSize);	69	+ task_args.push_back(opParam.recvDispls[rankIdx] * dataTypeSize);
71	}	70	}
72		71
73	- HcommCcuLaunchContextPod launch_context{};
74	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
75	- return HCCL_E_INTERNAL;
76	- }
77	CcuResult launchRet =	72	CcuResult launchRet =
78	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	73	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
79	return launchRet == CCU_SUCCESS ? HCCL_SUCCESS : HCCL_E_INTERNAL;	74	return launchRet == CCU_SUCCESS ? HCCL_SUCCESS : HCCL_E_INTERNAL;
80	}	75	}
81		76

Mexamples/ccu/ccu_all_to_all_vc/op_host/exec_op.h+1-1

	@@ -14,7 +14,7 @@
14	#include "alg_resource.h"	14	#include "alg_resource.h"
15		15
16	namespace ops_hccl_a2avc {	16	namespace ops_hccl_a2avc {
17	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	17	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
18	}	18	}
19		19
20	#endif	20	#endif

Mexamples/ccu/ccu_all_to_allv/op_host/alg_resource.cc+15-15

	@@ -22,33 +22,33 @@ namespace ops_hccl_a2av {
22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
23		23
24	static HcclResult AllocThreadAndChannelResource(	24	static HcclResult AllocThreadAndChannelResource(
25	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	25	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
26	uint32_t& kernelDie)	26	uint32_t& kernelDie)
27	{	27	{
28	ThreadHandle thread;	28	ThreadHandle thread;
29	constexpr uint32_t notifyNumOnMainThread = 0;	29	constexpr uint32_t notifyNumOnMainThread = 0;
30	RETURN_IF_HCCL_FAIL(	30	RETURN_IF_HCCL_FAIL(
31	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	31	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
32	resCtxHost.threads.push_back(thread);	32	resCtxHost.threads.push_back(thread);
33		33
34	- if (param_.rankSize == 1) {	34	+ if (opParam.rankSize == 1) {
35	return HCCL_SUCCESS;	35	return HCCL_SUCCESS;
36	}	36	}
37		37
38	- uint32_t channel_num = param_.rankSize - 1;	38	+ uint32_t channel_num = opParam.rankSize - 1;
39	kernelChannels.resize(channel_num);	39	kernelChannels.resize(channel_num);
40		40
41	uint32_t channelIndex = 0;	41	uint32_t channelIndex = 0;
42	bool hasKernelDie = false;	42	bool hasKernelDie = false;
43	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	43	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
44	- if (remoteRank == param_.myRank) {	44	+ if (remoteRank == opParam.myRank) {
45	continue;	45	continue;
46	}	46	}
47		47
48	uint32_t netLayer = 0, listSize = 0;	48	uint32_t netLayer = 0, listSize = 0;
49	CommLink* linkList = nullptr;	49	CommLink* linkList = nullptr;
50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(	50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(
51	- comm, netLayer, param_.myRank, remoteRank, &linkList,	51	+ comm, netLayer, opParam.myRank, remoteRank, &linkList,
52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息	52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息
53		53
54	HcclChannelDesc desc;	54	HcclChannelDesc desc;
	@@ -69,7 +69,7 @@ static HcclResult AllocThreadAndChannelResource(
69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
70	uint32_t channelDie = 0;	70	uint32_t channelDie = 0;
71	RETURN_IF_HCCL_FAIL(	71	RETURN_IF_HCCL_FAIL(
72	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	72	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
74	protocolExists = true;	74	protocolExists = true;
75	break;	75	break;
	@@ -88,7 +88,7 @@ static HcclResult AllocThreadAndChannelResource(
88	}	88	}
89		89
90	static HcclResult RegisterAllToAllVKernel(	90	static HcclResult RegisterAllToAllVKernel(
91	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	91	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
92	uint32_t kernelDie)	92	uint32_t kernelDie)
93	{	93	{
94	ccu_kernel_info kernelInfo;	94	ccu_kernel_info kernelInfo;
	@@ -98,8 +98,8 @@ static HcclResult RegisterAllToAllVKernel(
98	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllToAllVMesh1DKernel);	98	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllToAllVMesh1DKernel);
99		99
100	auto kernel_arg = std::make_shared<CcuKernelArgAllToAllVMesh1D>();	100	auto kernel_arg = std::make_shared<CcuKernelArgAllToAllVMesh1D>();
101	- kernel_arg->rankSize = param_.rankSize;	101	+ kernel_arg->rankSize = opParam.rankSize;
102	- kernel_arg->rankId = param_.myRank;	102	+ kernel_arg->rankId = opParam.myRank;
103	kernelInfo.setKernelArg(kernel_arg);	103	kernelInfo.setKernelArg(kernel_arg);
104		104
105	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	105	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -138,17 +138,17 @@ static HcclResult RegisterAllToAllVKernel(
138	return HCCL_SUCCESS;	138	return HCCL_SUCCESS;
139	}	139	}
140		140
141	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	141	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
142	{	142	{
143	std::vector<ChannelHandle> kernelChannels;	143	std::vector<ChannelHandle> kernelChannels;
144	uint32_t kernelDie = 0;	144	uint32_t kernelDie = 0;
145	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	145	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
146		146
147	- if (param_.rankSize == 1) {	147	+ if (opParam.rankSize == 1) {
148	return HCCL_SUCCESS;	148	return HCCL_SUCCESS;
149	}	149	}
150		150
151	- RETURN_IF_HCCL_FAIL(RegisterAllToAllVKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	151	+ RETURN_IF_HCCL_FAIL(RegisterAllToAllVKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
152	return HCCL_SUCCESS;	152	return HCCL_SUCCESS;
153	}	153	}
154		154

Mexamples/ccu/ccu_all_to_allv/op_host/alg_resource.h+1-1

	@@ -39,7 +39,7 @@ struct OpParam {
39	uint32_t rankSize = 0;	39	uint32_t rankSize = 0;
40	};	40	};
41		41
42	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	42	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
43		43
44	} // namespace ops_hccl_a2av	44	} // namespace ops_hccl_a2av
45		45

Mexamples/ccu/ccu_all_to_allv/op_host/exec_op.cc+16-21

	@@ -12,7 +12,6 @@
12	#include <vector>	12	#include <vector>
13	#include <ccu/hcomm/ccu_api_types.h>	13	#include <ccu/hcomm/ccu_api_types.h>
14	#include <hccl/hcomm_primitives.h>	14	#include <hccl/hcomm_primitives.h>
15	-#include <hcomm/hcomm_ccu_launch_api.h>
16	#include "ccu/hcomm/ccu_launch.h"	15	#include "ccu/hcomm/ccu_launch.h"
17	#include "ccu/hcomm/ccu_resource_api.h"	16	#include "ccu/hcomm/ccu_resource_api.h"
18	#include "alg_resource.h"	17	#include "alg_resource.h"
	@@ -32,50 +31,46 @@ static uint64_t GetBufferElements(const std::vector<uint64_t>& counts, const std
32	return elements;	31	return elements;
33	}	32	}
34		33
35	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	34	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
36	{	35	{
37	constexpr uint64_t dataTypeSize = sizeof(float);	36	constexpr uint64_t dataTypeSize = sizeof(float);
38	- const uint64_t inputSize = GetBufferElements(param_.sendCounts, param_.sendDispls) * dataTypeSize;	37	+ const uint64_t inputSize = GetBufferElements(opParam.sendCounts, opParam.sendDispls) * dataTypeSize;
39	- const uint64_t outputSize = GetBufferElements(param_.recvCounts, param_.recvDispls) * dataTypeSize;	38	+ const uint64_t outputSize = GetBufferElements(opParam.recvCounts, opParam.recvDispls) * dataTypeSize;
40	- if (param_.rankSize == 1) {	39	+ if (opParam.rankSize == 1) {
41	- if (param_.sendCounts[0] != param_.recvCounts[0]) {	40	+ if (opParam.sendCounts[0] != opParam.recvCounts[0]) {
42	return HCCL_E_PARA;	41	return HCCL_E_PARA;
43	}	42	}
44	- if (param_.sendCounts[0] == 0) {	43	+ if (opParam.sendCounts[0] == 0) {
45	return HCCL_SUCCESS;	44	return HCCL_SUCCESS;
46	}	45	}
47	- auto* src_ = static_cast<char>(param_.inputPtr) + param_.sendDispls[0] dataTypeSize;	46	+ auto* src_ = static_cast<char>(opParam.inputPtr) + opParam.sendDispls[0] dataTypeSize;
48	- auto* dst_ = static_cast<char>(param_.outputPtr) + param_.recvDispls[0] dataTypeSize;	47	+ auto* dst_ = static_cast<char>(opParam.outputPtr) + opParam.recvDispls[0] dataTypeSize;
49	return static_cast<HcclResult>(	48	return static_cast<HcclResult>(
50	- HcommLocalCopyOnThread(resCtx.threads[0], dst_, src_, param_.sendCounts[0] * dataTypeSize));	49	+ HcommLocalCopyOnThread(resCtx.threads[0], dst_, src_, opParam.sendCounts[0] * dataTypeSize));
51	}	50	}
52		51
53	uint64_t inputToken = 0;	52	uint64_t inputToken = 0;
54	uint64_t outputToken = 0;	53	uint64_t outputToken = 0;
55	- const uint64_t inputAddr = reinterpret_cast<uint64_t>(param_.inputPtr);	54	+ const uint64_t inputAddr = reinterpret_cast<uint64_t>(opParam.inputPtr);
56	- const uint64_t outputAddr = reinterpret_cast<uint64_t>(param_.outputPtr);	55	+ const uint64_t outputAddr = reinterpret_cast<uint64_t>(opParam.outputPtr);
57	if ((inputSize != 0 && asccomm_ccu_get_mem_token(inputAddr, inputSize, &inputToken) != CCU_SUCCESS) \|\|	56	if ((inputSize != 0 && asccomm_ccu_get_mem_token(inputAddr, inputSize, &inputToken) != CCU_SUCCESS) \|\|
58	(outputSize != 0 && asccomm_ccu_get_mem_token(outputAddr, outputSize, &outputToken) != CCU_SUCCESS)) {	57	(outputSize != 0 && asccomm_ccu_get_mem_token(outputAddr, outputSize, &outputToken) != CCU_SUCCESS)) {
59	return HCCL_E_INTERNAL;	58	return HCCL_E_INTERNAL;
60	}	59	}
61		60
62	std::vector<uint64_t> task_args = {inputAddr, outputAddr, inputToken, outputToken};	61	std::vector<uint64_t> task_args = {inputAddr, outputAddr, inputToken, outputToken};
63	- for (uint32_t rankIdx = 0; rankIdx < param_.rankSize; rankIdx++) {	62	+ for (uint32_t rankIdx = 0; rankIdx < opParam.rankSize; rankIdx++) {
64	- const uint64_t sendBytes = param_.sendCounts[rankIdx] * dataTypeSize;	63	+ const uint64_t sendBytes = opParam.sendCounts[rankIdx] * dataTypeSize;
65	if (sendBytes > MAX_TRANSFER_BYTES) {	64	if (sendBytes > MAX_TRANSFER_BYTES) {
66	return HCCL_E_NOT_SUPPORT;	65	return HCCL_E_NOT_SUPPORT;
67	}	66	}
68	task_args.push_back(sendBytes);	67	task_args.push_back(sendBytes);
69	- task_args.push_back(param_.sendDispls[rankIdx] * dataTypeSize);	68	+ task_args.push_back(opParam.sendDispls[rankIdx] * dataTypeSize);
70	- task_args.push_back(param_.recvDispls[rankIdx] * dataTypeSize);	69	+ task_args.push_back(opParam.recvDispls[rankIdx] * dataTypeSize);
71	}	70	}
72		71
73	- HcommCcuLaunchContextPod launch_context{};
74	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
75	- return HCCL_E_INTERNAL;
76	- }
77	CcuResult launchRet =	72	CcuResult launchRet =
78	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	73	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
79	return launchRet == CCU_SUCCESS ? HCCL_SUCCESS : HCCL_E_INTERNAL;	74	return launchRet == CCU_SUCCESS ? HCCL_SUCCESS : HCCL_E_INTERNAL;
80	}	75	}
81		76

Mexamples/ccu/ccu_all_to_allv/op_host/exec_op.h+1-1

	@@ -14,7 +14,7 @@
14	#include "alg_resource.h"	14	#include "alg_resource.h"
15		15
16	namespace ops_hccl_a2av {	16	namespace ops_hccl_a2av {
17	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	17	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
18	}	18	}
19		19
20	#endif	20	#endif

Mexamples/ccu/ccu_allgather/op_host/alg_resource.cc+15-15

	@@ -25,33 +25,33 @@ namespace ops_hccl_ag {
25	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	25	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
26		26
27	static HcclResult AllocThreadAndChannelResource(	27	static HcclResult AllocThreadAndChannelResource(
28	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	28	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
29	uint32_t& kernelDie)	29	uint32_t& kernelDie)
30	{	30	{
31	ThreadHandle thread;	31	ThreadHandle thread;
32	constexpr uint32_t notifyNumOnMainThread = 0;	32	constexpr uint32_t notifyNumOnMainThread = 0;
33	RETURN_IF_HCCL_FAIL(	33	RETURN_IF_HCCL_FAIL(
34	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	34	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
35	resCtxHost.threads.push_back(thread);	35	resCtxHost.threads.push_back(thread);
36		36
37	- if (param_.rankSize == 1) {	37	+ if (opParam.rankSize == 1) {
38	return HCCL_SUCCESS;	38	return HCCL_SUCCESS;
39	}	39	}
40		40
41	- uint32_t channel_num = param_.rankSize - 1;	41	+ uint32_t channel_num = opParam.rankSize - 1;
42	kernelChannels.resize(channel_num);	42	kernelChannels.resize(channel_num);
43		43
44	uint32_t channelIndex = 0;	44	uint32_t channelIndex = 0;
45	bool hasKernelDie = false;	45	bool hasKernelDie = false;
46	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	46	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
47	- if (remoteRank == param_.myRank) {	47	+ if (remoteRank == opParam.myRank) {
48	continue;	48	continue;
49	}	49	}
50		50
51	uint32_t netLayer = 0;	51	uint32_t netLayer = 0;
52	uint32_t listSize = 0;	52	uint32_t listSize = 0;
53	CommLink* linkList = nullptr;	53	CommLink* linkList = nullptr;
54	- RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(comm, netLayer, param_.myRank, remoteRank, &linkList, &listSize));	54	+ RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(comm, netLayer, opParam.myRank, remoteRank, &linkList, &listSize));
55		55
56	HcclChannelDesc desc;	56	HcclChannelDesc desc;
57	RETURN_IF_HCCL_FAIL(HcclChannelDescInit(&desc, 1));	57	RETURN_IF_HCCL_FAIL(HcclChannelDescInit(&desc, 1));
	@@ -71,7 +71,7 @@ static HcclResult AllocThreadAndChannelResource(
71	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	71	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
72	uint32_t channelDie = 0;	72	uint32_t channelDie = 0;
73	RETURN_IF_HCCL_FAIL(	73	RETURN_IF_HCCL_FAIL(
74	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	74	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
75	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	75	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
76	protocolExists = true;	76	protocolExists = true;
77	break;	77	break;
	@@ -89,7 +89,7 @@ static HcclResult AllocThreadAndChannelResource(
89	}	89	}
90		90
91	static HcclResult RegisterAllGatherKernel(	91	static HcclResult RegisterAllGatherKernel(
92	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	92	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
93	uint32_t kernelDie)	93	uint32_t kernelDie)
94	{	94	{
95	ccu_kernel_info kernelInfo;	95	ccu_kernel_info kernelInfo;
	@@ -100,8 +100,8 @@ static HcclResult RegisterAllGatherKernel(
100	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllGatherMesh1DMem2MemKernel);	100	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllGatherMesh1DMem2MemKernel);
101		101
102	auto kernel_arg = std::make_shared<CcuKernelArgAllGatherMesh1DMem2Mem>();	102	auto kernel_arg = std::make_shared<CcuKernelArgAllGatherMesh1DMem2Mem>();
103	- kernel_arg->rankSize = param_.rankSize;	103	+ kernel_arg->rankSize = opParam.rankSize;
104	- kernel_arg->rankId = param_.myRank;	104	+ kernel_arg->rankId = opParam.myRank;
105	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	105	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
106	kernel_arg->channels[i] = kernelChannels[i];	106	kernel_arg->channels[i] = kernelChannels[i];
107	}	107	}
	@@ -138,17 +138,17 @@ static HcclResult RegisterAllGatherKernel(
138	return HCCL_SUCCESS;	138	return HCCL_SUCCESS;
139	}	139	}
140		140
141	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	141	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
142	{	142	{
143	std::vector<ChannelHandle> kernelChannels;	143	std::vector<ChannelHandle> kernelChannels;
144	uint32_t kernelDie = 0;	144	uint32_t kernelDie = 0;
145	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	145	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
146		146
147	- if (param_.rankSize == 1) {	147	+ if (opParam.rankSize == 1) {
148	return HCCL_SUCCESS;	148	return HCCL_SUCCESS;
149	}	149	}
150		150
151	- RETURN_IF_HCCL_FAIL(RegisterAllGatherKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	151	+ RETURN_IF_HCCL_FAIL(RegisterAllGatherKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
152	return HCCL_SUCCESS;	152	return HCCL_SUCCESS;
153	}	153	}
154		154

Mexamples/ccu/ccu_allgather/op_host/alg_resource.h+1-1

	@@ -37,7 +37,7 @@ struct OpParam {
37	uint32_t rankSize = 0;	37	uint32_t rankSize = 0;
38	};	38	};
39		39
40	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	40	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
41		41
42	} // namespace ops_hccl_ag	42	} // namespace ops_hccl_ag
43		43

Mexamples/ccu/ccu_allgather/op_host/exec_op.cc+16-21

	@@ -14,7 +14,6 @@
14		14
15	#include <ccu/hcomm/ccu_api_types.h>	15	#include <ccu/hcomm/ccu_api_types.h>
16	#include <hccl/hcomm_primitives.h>	16	#include <hccl/hcomm_primitives.h>
17	-#include <hcomm/hcomm_ccu_launch_api.h>
18		17
19	#include "alg_resource.h"	18	#include "alg_resource.h"
20	#include "ccu/hcomm/ccu_launch.h"	19	#include "ccu/hcomm/ccu_launch.h"
	@@ -34,13 +33,13 @@ static HcclResult GetDataTypeSize(HcclDataType data_type_, uint32_t& dataTypeSiz
34	}	33	}
35		34
36	static HcclResult LaunchCcuKernelSlice(	35	static HcclResult LaunchCcuKernelSlice(
37	- const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t token, uint64_t dataSize,	36	+ const OpParam& opParam, const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t token,
38	- uint32_t myRank, uint64_t sliceCount, uint64_t dataTypeSize)	37	+ uint64_t dataSize, uint64_t sliceCount, uint64_t dataTypeSize)
39	{	38	{
40	uint64_t sliceSize = sliceCount * dataTypeSize;	39	uint64_t sliceSize = sliceCount * dataTypeSize;
41		40
42	uint64_t currentRankSliceInputOffset = 0;	41	uint64_t currentRankSliceInputOffset = 0;
43	- uint64_t currentRankSliceOutputOffset = dataSize * myRank;	42	+ uint64_t currentRankSliceOutputOffset = dataSize * opParam.myRank;
44		43
45	loop_group_config config{};	44	loop_group_config config{};
46	config.ms_interleave = ccu_ms_interleave;	45	config.ms_interleave = ccu_ms_interleave;
	@@ -61,12 +60,8 @@ static HcclResult LaunchCcuKernelSlice(
61	goSize.residual,	60	goSize.residual,
62	};	61	};
63		62
64	- HcommCcuLaunchContextPod launch_context{};
65	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
66	- return HCCL_E_INTERNAL;
67	- }
68	CcuResult launchRet =	63	CcuResult launchRet =
69	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	64	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
70	if (launchRet != CCU_SUCCESS) {	65	if (launchRet != CCU_SUCCESS) {
71	return HCCL_E_INTERNAL;	66	return HCCL_E_INTERNAL;
72	}	67	}
	@@ -74,20 +69,20 @@ static HcclResult LaunchCcuKernelSlice(
74	return HCCL_SUCCESS;	69	return HCCL_SUCCESS;
75	}	70	}
76		71
77	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	72	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
78	{	73	{
79	uint32_t dataTypeSize = 0;	74	uint32_t dataTypeSize = 0;
80	- RETURN_IF_HCCL_FAIL(GetDataTypeSize(param_.data_type_, dataTypeSize));	75	+ RETURN_IF_HCCL_FAIL(GetDataTypeSize(opParam.data_type_, dataTypeSize));
81	- uint64_t dataSize = param_.count_ * dataTypeSize;	76	+ uint64_t dataSize = opParam.count_ * dataTypeSize;
82	- uint64_t count_ = param_.count_;	77	+ uint64_t count_ = opParam.count_;
83		78
84	if (count_ == 0) {	79	if (count_ == 0) {
85	return HCCL_SUCCESS;	80	return HCCL_SUCCESS;
86	}	81	}
87		82
88	- if (param_.rankSize == 1) {	83	+ if (opParam.rankSize == 1) {
89	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(	84	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(
90	- HcommLocalCopyOnThread(resCtx.threads[0], param_.outputPtr, param_.inputPtr, dataSize)));	85	+ HcommLocalCopyOnThread(resCtx.threads[0], opParam.outputPtr, opParam.inputPtr, dataSize)));
91	return HCCL_SUCCESS;	86	return HCCL_SUCCESS;
92	}	87	}
93		88
	@@ -96,11 +91,11 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
96	uint64_t processedDataCount = 0;	91	uint64_t processedDataCount = 0;
97		92
98	uint64_t token = 0;	93	uint64_t token = 0;
99	- uint64_t baseInputAddr = reinterpret_cast<uint64_t>(param_.inputPtr);	94	+ uint64_t baseInputAddr = reinterpret_cast<uint64_t>(opParam.inputPtr);
100	- uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(param_.outputPtr);	95	+ uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(opParam.outputPtr);
101	- if (param_.inputPtr != nullptr) {	96	+ if (opParam.inputPtr != nullptr) {
102	asccomm_ccu_get_mem_token(baseInputAddr, static_cast<uint64_t>(dataSize), &token);	97	asccomm_ccu_get_mem_token(baseInputAddr, static_cast<uint64_t>(dataSize), &token);
103	- } else if (param_.outputPtr != nullptr) {	98	+ } else if (opParam.outputPtr != nullptr) {
104	asccomm_ccu_get_mem_token(baseOutputAddr, static_cast<uint64_t>(dataSize), &token);	99	asccomm_ccu_get_mem_token(baseOutputAddr, static_cast<uint64_t>(dataSize), &token);
105	}	100	}
106		101
	@@ -109,8 +104,8 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
109	uint64_t inputAddr = baseInputAddr + processedDataCount * dataTypeSize;	104	uint64_t inputAddr = baseInputAddr + processedDataCount * dataTypeSize;
110	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;	105	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;
111		106
112	- RETURN_IF_HCCL_FAIL(LaunchCcuKernelSlice(	107	+ RETURN_IF_HCCL_FAIL(
113	- resCtx, inputAddr, outputAddr, token, dataSize, param_.myRank, sliceCount, dataTypeSize));	108	+ LaunchCcuKernelSlice(opParam, resCtx, inputAddr, outputAddr, token, dataSize, sliceCount, dataTypeSize));
114		109
115	processedDataCount += sliceCount;	110	processedDataCount += sliceCount;
116	}	111	}

Mexamples/ccu/ccu_allgather/op_host/exec_op.h+1-1

	@@ -15,7 +15,7 @@
15		15
16	namespace ops_hccl_ag {	16	namespace ops_hccl_ag {
17		17
18	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	18	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
19		19
20	} // namespace ops_hccl_ag	20	} // namespace ops_hccl_ag
21		21

Mexamples/ccu/ccu_allgather_mesh1d_mem2mem/op_host/exec_op.asc+1-4

	@@ -24,7 +24,6 @@
24	#include <acl/acl.h>	24	#include <acl/acl.h>
25	#include <hcomm/hcomm_types.h>	25	#include <hcomm/hcomm_types.h>
26	#include <hcomm/hcomm_primitives.h>	26	#include <hcomm/hcomm_primitives.h>
27	-#include <hcomm/hcomm_ccu_launch_api.h>
28		27
29	#include "alg_resource.h"	28	#include "alg_resource.h"
30	#include "ccu_kernel.h"	29	#include "ccu_kernel.h"
	@@ -151,10 +150,8 @@ HcclResult AllGatherMesh1D(
151	task_args[TASK_GO_PARALLEL_PARAM] = goSize[2];	150	task_args[TASK_GO_PARALLEL_PARAM] = goSize[2];
152	task_args[TASK_GO_RESIDUAL] = goSize[3];	151	task_args[TASK_GO_RESIDUAL] = goSize[3];
153		152
154	- HcommCcuLaunchContextPod launch_context{};
155	- RETURN_IF_HCCL_FAIL(HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context));
156	RETURN_IF_HCCL_FAIL(	153	RETURN_IF_HCCL_FAIL(
157	- ConvertCcuToHccl(asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args, TASK_ARG_NUM)));	154	+ ConvertCcuToHccl(asccomm_ccu_kernel_launch(stream, resCtx.ccuKernels[0], task_args, TASK_ARG_NUM)));
158	return HCCL_SUCCESS;	155	return HCCL_SUCCESS;
159	}	156	}
160		157

Mexamples/ccu/ccu_allgather_nhr1d_mem2mem/op_host/exec_op.asc+1-4

	@@ -25,7 +25,6 @@
25	#include <acl/acl.h>	25	#include <acl/acl.h>
26	#include <hcomm/hcomm_types.h>	26	#include <hcomm/hcomm_types.h>
27	#include <hcomm/hcomm_primitives.h>	27	#include <hcomm/hcomm_primitives.h>
28	-#include <hcomm/hcomm_ccu_launch_api.h>
29		28
30	#include <ccu/hcomm/ccu_launch.h>	29	#include <ccu/hcomm/ccu_launch.h>
31	#include <ccu/hcomm/ccu_resource_api.h>	30	#include <ccu/hcomm/ccu_resource_api.h>
	@@ -153,10 +152,8 @@ HcclResult AllGatherNhr1D(
153	task_args[TASK_GO_PARALLEL_PARAM] = goSize[2];	152	task_args[TASK_GO_PARALLEL_PARAM] = goSize[2];
154	task_args[TASK_GO_RESIDUAL] = goSize[3];	153	task_args[TASK_GO_RESIDUAL] = goSize[3];
155		154
156	- HcommCcuLaunchContextPod launch_context{};
157	- RETURN_IF_HCCL_FAIL(HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context));
158	RETURN_IF_HCCL_FAIL(	155	RETURN_IF_HCCL_FAIL(
159	- ConvertCcuToHccl(asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args, TASK_ARG_NUM)));	156	+ ConvertCcuToHccl(asccomm_ccu_kernel_launch(stream, resCtx.ccuKernels[0], task_args, TASK_ARG_NUM)));
160	return HCCL_SUCCESS;	157	return HCCL_SUCCESS;
161	}	158	}
162		159

Mexamples/ccu/ccu_allreduce/op_host/alg_resource.cc+15-15

	@@ -23,33 +23,33 @@ namespace ops_hccl_ar {
23	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	23	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
24		24
25	static HcclResult AllocThreadAndChannelResource(	25	static HcclResult AllocThreadAndChannelResource(
26	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	26	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
27	uint32_t& kernelDie)	27	uint32_t& kernelDie)
28	{	28	{
29	ThreadHandle thread;	29	ThreadHandle thread;
30	constexpr uint32_t notifyNumOnMainThread = 0;	30	constexpr uint32_t notifyNumOnMainThread = 0;
31	RETURN_IF_HCCL_FAIL(	31	RETURN_IF_HCCL_FAIL(
32	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	32	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
33	resCtxHost.threads.push_back(thread);	33	resCtxHost.threads.push_back(thread);
34		34
35	- if (param_.rankSize == 1) {	35	+ if (opParam.rankSize == 1) {
36	return HCCL_SUCCESS;	36	return HCCL_SUCCESS;
37	}	37	}
38		38
39	- uint32_t channel_num = param_.rankSize - 1;	39	+ uint32_t channel_num = opParam.rankSize - 1;
40	kernelChannels.resize(channel_num);	40	kernelChannels.resize(channel_num);
41		41
42	uint32_t channelIndex = 0;	42	uint32_t channelIndex = 0;
43	bool hasKernelDie = false;	43	bool hasKernelDie = false;
44	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	44	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
45	- if (remoteRank == param_.myRank) {	45	+ if (remoteRank == opParam.myRank) {
46	continue;	46	continue;
47	}	47	}
48		48
49	uint32_t netLayer = 0, listSize = 0;	49	uint32_t netLayer = 0, listSize = 0;
50	CommLink* linkList = nullptr;	50	CommLink* linkList = nullptr;
51	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(	51	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(
52	- comm, netLayer, param_.myRank, remoteRank, &linkList,	52	+ comm, netLayer, opParam.myRank, remoteRank, &linkList,
53	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息	53	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息
54		54
55	HcclChannelDesc desc;	55	HcclChannelDesc desc;
	@@ -70,7 +70,7 @@ static HcclResult AllocThreadAndChannelResource(
70	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	70	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
71	uint32_t channelDie = 0;	71	uint32_t channelDie = 0;
72	RETURN_IF_HCCL_FAIL(	72	RETURN_IF_HCCL_FAIL(
73	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	73	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
74	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	74	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
75	protocolExists = true;	75	protocolExists = true;
76	break;	76	break;
	@@ -89,7 +89,7 @@ static HcclResult AllocThreadAndChannelResource(
89	}	89	}
90		90
91	static HcclResult RegisterAllReduceKernel(	91	static HcclResult RegisterAllReduceKernel(
92	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	92	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
93	uint32_t kernelDie)	93	uint32_t kernelDie)
94	{	94	{
95	ccu_kernel_info kernelInfo;	95	ccu_kernel_info kernelInfo;
	@@ -100,8 +100,8 @@ static HcclResult RegisterAllReduceKernel(
100	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllReduceMesh1DMem2MemKernel);	100	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuAllReduceMesh1DMem2MemKernel);
101		101
102	auto kernel_arg = std::make_shared<CcuKernelArgAllReduceMesh1DMem2Mem>();	102	auto kernel_arg = std::make_shared<CcuKernelArgAllReduceMesh1DMem2Mem>();
103	- kernel_arg->rankSize = param_.rankSize;	103	+ kernel_arg->rankSize = opParam.rankSize;
104	- kernel_arg->rankId = param_.myRank;	104	+ kernel_arg->rankId = opParam.myRank;
105	kernelInfo.setKernelArg(kernel_arg);	105	kernelInfo.setKernelArg(kernel_arg);
106		106
107	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	107	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -140,17 +140,17 @@ static HcclResult RegisterAllReduceKernel(
140	return HCCL_SUCCESS;	140	return HCCL_SUCCESS;
141	}	141	}
142		142
143	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	143	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
144	{	144	{
145	std::vector<ChannelHandle> kernelChannels;	145	std::vector<ChannelHandle> kernelChannels;
146	uint32_t kernelDie = 0;	146	uint32_t kernelDie = 0;
147	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	147	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
148		148
149	- if (param_.rankSize == 1) {	149	+ if (opParam.rankSize == 1) {
150	return HCCL_SUCCESS;	150	return HCCL_SUCCESS;
151	}	151	}
152		152
153	- RETURN_IF_HCCL_FAIL(RegisterAllReduceKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	153	+ RETURN_IF_HCCL_FAIL(RegisterAllReduceKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
154	return HCCL_SUCCESS;	154	return HCCL_SUCCESS;
155	}	155	}
156		156

Mexamples/ccu/ccu_allreduce/op_host/alg_resource.h+1-1

	@@ -36,7 +36,7 @@ struct OpParam {
36	uint32_t rankSize = 0;	36	uint32_t rankSize = 0;
37	};	37	};
38		38
39	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	39	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
40		40
41	} // namespace ops_hccl_ar	41	} // namespace ops_hccl_ar
42		42

Mexamples/ccu/ccu_allreduce/op_host/exec_op.cc+14-18

	@@ -12,7 +12,6 @@
12	#include <algorithm>	12	#include <algorithm>
13	#include <ccu/hcomm/ccu_api_types.h>	13	#include <ccu/hcomm/ccu_api_types.h>
14	#include <hccl/hcomm_primitives.h>	14	#include <hccl/hcomm_primitives.h>
15	-#include <hcomm/hcomm_ccu_launch_api.h>
16	#include "ccu/hcomm/ccu_launch.h"	15	#include "ccu/hcomm/ccu_launch.h"
17	#include "ccu/hcomm/ccu_resource_api.h"	16	#include "ccu/hcomm/ccu_resource_api.h"
18	#include "alg_resource.h"	17	#include "alg_resource.h"
	@@ -23,8 +22,8 @@ namespace ops_hccl_ar {
23	constexpr uint64_t UB_MAX_DATA_SIZE = 256 * 1024 * 1024; // UB 协议单次传输最大字节数	22	constexpr uint64_t UB_MAX_DATA_SIZE = 256 * 1024 * 1024; // UB 协议单次传输最大字节数
24		23
25	static HcclResult LaunchCcuKernelSlice(	24	static HcclResult LaunchCcuKernelSlice(
26	- const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t token, uint64_t sliceCount,	25	+ const OpParam& opParam, const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t token,
27	- uint64_t dataTypeSize)	26	+ uint64_t sliceCount, uint64_t dataTypeSize)
28	{	27	{
29	uint64_t sliceSize = sliceCount * dataTypeSize;	28	uint64_t sliceSize = sliceCount * dataTypeSize;
30		29
	@@ -35,12 +34,8 @@ static HcclResult LaunchCcuKernelSlice(
35	sliceSize,	34	sliceSize,
36	};	35	};
37		36
38	- HcommCcuLaunchContextPod launch_context{};
39	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
40	- return HCCL_E_INTERNAL;
41	- }
42	CcuResult launchRet =	37	CcuResult launchRet =
43	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	38	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
44	if (launchRet != CCU_SUCCESS) {	39	if (launchRet != CCU_SUCCESS) {
45	return HCCL_E_INTERNAL;	40	return HCCL_E_INTERNAL;
46	}	41	}
	@@ -48,19 +43,19 @@ static HcclResult LaunchCcuKernelSlice(
48	return HCCL_SUCCESS;	43	return HCCL_SUCCESS;
49	}	44	}
50		45
51	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	46	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
52	{	47	{
53	constexpr uint64_t dataTypeSize = sizeof(float);	48	constexpr uint64_t dataTypeSize = sizeof(float);
54	- uint64_t dataSize = param_.count_ * dataTypeSize;	49	+ uint64_t dataSize = opParam.count_ * dataTypeSize;
55	- uint64_t count_ = param_.count_;	50	+ uint64_t count_ = opParam.count_;
56		51
57	if (count_ == 0) { // 数据量为 0，直接返回	52	if (count_ == 0) { // 数据量为 0，直接返回
58	return HcclResult::HCCL_SUCCESS;	53	return HcclResult::HCCL_SUCCESS;
59	}	54	}
60		55
61	- if (param_.rankSize == 1) { // 单卡场景直接本地拷贝	56	+ if (opParam.rankSize == 1) { // 单卡场景直接本地拷贝
62	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(	57	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(
63	- HcommLocalCopyOnThread(resCtx.threads[0], param_.outputPtr, param_.inputPtr, dataSize)));	58	+ HcommLocalCopyOnThread(resCtx.threads[0], opParam.outputPtr, opParam.inputPtr, dataSize)));
64	return HCCL_SUCCESS;	59	return HCCL_SUCCESS;
65	}	60	}
66		61
	@@ -71,11 +66,11 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
71	uint64_t processedDataCount = 0;	66	uint64_t processedDataCount = 0;
72		67
73	uint64_t token = 0;	68	uint64_t token = 0;
74	- uint64_t baseInputAddr = reinterpret_cast<uint64_t>(param_.inputPtr);	69	+ uint64_t baseInputAddr = reinterpret_cast<uint64_t>(opParam.inputPtr);
75	- uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(param_.outputPtr);	70	+ uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(opParam.outputPtr);
76	- if (param_.inputPtr != nullptr) {	71	+ if (opParam.inputPtr != nullptr) {
77	asccomm_ccu_get_mem_token(baseInputAddr, static_cast<uint64_t>(dataSize), &token);	72	asccomm_ccu_get_mem_token(baseInputAddr, static_cast<uint64_t>(dataSize), &token);
78	- } else if (param_.outputPtr != nullptr) {	73	+ } else if (opParam.outputPtr != nullptr) {
79	asccomm_ccu_get_mem_token(baseOutputAddr, static_cast<uint64_t>(dataSize), &token);	74	asccomm_ccu_get_mem_token(baseOutputAddr, static_cast<uint64_t>(dataSize), &token);
80	}	75	}
81		76
	@@ -84,7 +79,8 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
84	uint64_t inputAddr = baseInputAddr + processedDataCount * dataTypeSize;	79	uint64_t inputAddr = baseInputAddr + processedDataCount * dataTypeSize;
85	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;	80	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;
86		81
87	- RETURN_IF_HCCL_FAIL(LaunchCcuKernelSlice(resCtx, inputAddr, outputAddr, token, sliceCount, dataTypeSize));	82	+ RETURN_IF_HCCL_FAIL(
		83	+ LaunchCcuKernelSlice(opParam, resCtx, inputAddr, outputAddr, token, sliceCount, dataTypeSize));
88		84
89	processedDataCount += sliceCount;	85	processedDataCount += sliceCount;
90	}	86	}

Mexamples/ccu/ccu_allreduce/op_host/exec_op.h+1-1

	@@ -15,7 +15,7 @@
15		15
16	namespace ops_hccl_ar {	16	namespace ops_hccl_ar {
17		17
18	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	18	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
19		19
20	} // namespace ops_hccl_ar	20	} // namespace ops_hccl_ar
21		21

Mexamples/ccu/ccu_broadcast/op_host/alg_resource.cc+16-16

	@@ -22,33 +22,33 @@ namespace ops_hccl_bcast {
22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
23		23
24	static HcclResult AllocThreadAndChannelResource(	24	static HcclResult AllocThreadAndChannelResource(
25	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	25	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
26	uint32_t& kernelDie)	26	uint32_t& kernelDie)
27	{	27	{
28	ThreadHandle thread;	28	ThreadHandle thread;
29	constexpr uint32_t notifyNumOnMainThread = 0;	29	constexpr uint32_t notifyNumOnMainThread = 0;
30	RETURN_IF_HCCL_FAIL(	30	RETURN_IF_HCCL_FAIL(
31	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	31	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
32	resCtxHost.threads.push_back(thread);	32	resCtxHost.threads.push_back(thread);
33		33
34	- if (param_.rankSize == 1) {	34	+ if (opParam.rankSize == 1) {
35	return HCCL_SUCCESS;	35	return HCCL_SUCCESS;
36	}	36	}
37		37
38	- uint32_t channel_num = param_.rankSize - 1;	38	+ uint32_t channel_num = opParam.rankSize - 1;
39	kernelChannels.resize(channel_num);	39	kernelChannels.resize(channel_num);
40		40
41	uint32_t channelIndex = 0;	41	uint32_t channelIndex = 0;
42	bool hasKernelDie = false;	42	bool hasKernelDie = false;
43	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	43	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
44	- if (remoteRank == param_.myRank) {	44	+ if (remoteRank == opParam.myRank) {
45	continue;	45	continue;
46	}	46	}
47		47
48	uint32_t netLayer = 0, listSize = 0;	48	uint32_t netLayer = 0, listSize = 0;
49	CommLink* linkList = nullptr;	49	CommLink* linkList = nullptr;
50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(	50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(
51	- comm, netLayer, param_.myRank, remoteRank, &linkList,	51	+ comm, netLayer, opParam.myRank, remoteRank, &linkList,
52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息	52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息
53		53
54	HcclChannelDesc desc;	54	HcclChannelDesc desc;
	@@ -69,7 +69,7 @@ static HcclResult AllocThreadAndChannelResource(
69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
70	uint32_t channelDie = 0;	70	uint32_t channelDie = 0;
71	RETURN_IF_HCCL_FAIL(	71	RETURN_IF_HCCL_FAIL(
72	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	72	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
74	protocolExists = true;	74	protocolExists = true;
75	break;	75	break;
	@@ -88,7 +88,7 @@ static HcclResult AllocThreadAndChannelResource(
88	}	88	}
89		89
90	static HcclResult RegisterBroadcastKernel(	90	static HcclResult RegisterBroadcastKernel(
91	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	91	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
92	uint32_t kernelDie)	92	uint32_t kernelDie)
93	{	93	{
94	ccu_kernel_info kernelInfo;	94	ccu_kernel_info kernelInfo;
	@@ -98,9 +98,9 @@ static HcclResult RegisterBroadcastKernel(
98	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuBroadcastMesh1DKernel);	98	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuBroadcastMesh1DKernel);
99		99
100	auto kernel_arg = std::make_shared<CcuKernelArgBroadcastMesh1D>();	100	auto kernel_arg = std::make_shared<CcuKernelArgBroadcastMesh1D>();
101	- kernel_arg->rankSize = param_.rankSize;	101	+ kernel_arg->rankSize = opParam.rankSize;
102	- kernel_arg->rankId = param_.myRank;	102	+ kernel_arg->rankId = opParam.myRank;
103	- kernel_arg->rootId = param_.rootRank;	103	+ kernel_arg->rootId = opParam.rootRank;
104	kernelInfo.setKernelArg(kernel_arg);	104	kernelInfo.setKernelArg(kernel_arg);
105		105
106	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	106	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -139,17 +139,17 @@ static HcclResult RegisterBroadcastKernel(
139	return HCCL_SUCCESS;	139	return HCCL_SUCCESS;
140	}	140	}
141		141
142	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	142	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
143	{	143	{
144	std::vector<ChannelHandle> kernelChannels;	144	std::vector<ChannelHandle> kernelChannels;
145	uint32_t kernelDie = 0;	145	uint32_t kernelDie = 0;
146	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	146	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
147		147
148	- if (param_.rankSize == 1) {	148	+ if (opParam.rankSize == 1) {
149	return HCCL_SUCCESS;	149	return HCCL_SUCCESS;
150	}	150	}
151		151
152	- RETURN_IF_HCCL_FAIL(RegisterBroadcastKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	152	+ RETURN_IF_HCCL_FAIL(RegisterBroadcastKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
153	return HCCL_SUCCESS;	153	return HCCL_SUCCESS;
154	}	154	}
155		155

Mexamples/ccu/ccu_broadcast/op_host/alg_resource.h+1-1

	@@ -36,7 +36,7 @@ struct OpParam {
36	uint32_t rankSize = 0;	36	uint32_t rankSize = 0;
37	};	37	};
38		38
39	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	39	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
40		40
41	} // namespace ops_hccl_bcast	41	} // namespace ops_hccl_bcast
42		42

Mexamples/ccu/ccu_broadcast/op_host/exec_op.cc+5-10

	@@ -10,7 +10,6 @@
10		10
11	#include <vector>	11	#include <vector>
12	#include <ccu/hcomm/ccu_api_types.h>	12	#include <ccu/hcomm/ccu_api_types.h>
13	-#include <hcomm/hcomm_ccu_launch_api.h>
14	#include "ccu/hcomm/ccu_launch.h"	13	#include "ccu/hcomm/ccu_launch.h"
15	#include "ccu/hcomm/ccu_resource_api.h"	14	#include "ccu/hcomm/ccu_resource_api.h"
16	#include "alg_resource.h"	15	#include "alg_resource.h"
	@@ -19,28 +18,24 @@ namespace ops_hccl_bcast {
19		18
20	constexpr uint64_t MAX_TRANSFER_BYTES = 256ULL * 1024 * 1024;	19	constexpr uint64_t MAX_TRANSFER_BYTES = 256ULL * 1024 * 1024;
21		20
22	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	21	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
23	{	22	{
24	- const uint64_t dataSize = param_.count_ * sizeof(float);	23	+ const uint64_t dataSize = opParam.count_ * sizeof(float);
25	- if (dataSize == 0 \|\| param_.rankSize == 1) {	24	+ if (dataSize == 0 \|\| opParam.rankSize == 1) {
26	return HCCL_SUCCESS;	25	return HCCL_SUCCESS;
27	}	26	}
28	if (dataSize > MAX_TRANSFER_BYTES) {	27	if (dataSize > MAX_TRANSFER_BYTES) {
29	return HCCL_E_NOT_SUPPORT;	28	return HCCL_E_NOT_SUPPORT;
30	}	29	}
31		30
32	- const uint64_t bufferAddr = reinterpret_cast<uint64_t>(param_.buffer);	31	+ const uint64_t bufferAddr = reinterpret_cast<uint64_t>(opParam.buffer);
33	uint64_t token = 0;	32	uint64_t token = 0;
34	if (asccomm_ccu_get_mem_token(bufferAddr, dataSize, &token) != CCU_SUCCESS) {	33	if (asccomm_ccu_get_mem_token(bufferAddr, dataSize, &token) != CCU_SUCCESS) {
35	return HCCL_E_INTERNAL;	34	return HCCL_E_INTERNAL;
36	}	35	}
37	std::vector<uint64_t> task_args = {bufferAddr, token, dataSize};	36	std::vector<uint64_t> task_args = {bufferAddr, token, dataSize};
38	- HcommCcuLaunchContextPod launch_context{};
39	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
40	- return HCCL_E_INTERNAL;
41	- }
42	CcuResult launchRet =	37	CcuResult launchRet =
43	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	38	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
44	return launchRet == CCU_SUCCESS ? HCCL_SUCCESS : HCCL_E_INTERNAL;	39	return launchRet == CCU_SUCCESS ? HCCL_SUCCESS : HCCL_E_INTERNAL;
45	}	40	}
46		41

Mexamples/ccu/ccu_broadcast/op_host/exec_op.h+1-1

	@@ -14,7 +14,7 @@
14	#include "alg_resource.h"	14	#include "alg_resource.h"
15		15
16	namespace ops_hccl_bcast {	16	namespace ops_hccl_bcast {
17	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	17	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
18	}	18	}
19		19
20	#endif	20	#endif

Mexamples/ccu/ccu_direct/01_allgather/CMakeLists.txt+5-5

	@@ -30,12 +30,10 @@ target_include_directories(demo PRIVATE
30	# CANN Toolkit头文件	30	# CANN Toolkit头文件
31	${ASCEND_CANN_PACKAGE_PATH}/include	31	${ASCEND_CANN_PACKAGE_PATH}/include
32	${ASCEND_CANN_PACKAGE_PATH}/include/hccl	32	${ASCEND_CANN_PACKAGE_PATH}/include/hccl
33	- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm	33	+ ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api
34	- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm/ccu
35	- ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api/ccu
36	${ASCEND_CANN_PACKAGE_PATH}/pkg_inc	34	${ASCEND_CANN_PACKAGE_PATH}/pkg_inc
37	- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/hcomm	35	+ # 样例公共辅助头（ccu_instance.h）
38	- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/hcomm/ccu	36	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../common
39	)	37	)
40		38
41	target_compile_features(demo PRIVATE cxx_std_14)	39	target_compile_features(demo PRIVATE cxx_std_14)
	@@ -59,6 +57,7 @@ target_link_options(demo PRIVATE
59	-Wl,-z,relro	57	-Wl,-z,relro
60	-Wl,-z,now	58	-Wl,-z,now
61	-Wl,-z,noexecstack	59	-Wl,-z,noexecstack
		60	+ -Wl,-rpath,${ASCEND_CANN_PACKAGE_PATH}/lib64
62	)	61	)
63	target_link_directories(demo PRIVATE	62	target_link_directories(demo PRIVATE
64	${ASCEND_CANN_PACKAGE_PATH}/lib64	63	${ASCEND_CANN_PACKAGE_PATH}/lib64
	@@ -66,6 +65,7 @@ target_link_directories(demo PRIVATE
66	target_link_libraries(demo PRIVATE	65	target_link_libraries(demo PRIVATE
67	hcomm	66	hcomm
68	asccomm_ccu	67	asccomm_ccu
		68	+ asccomm_ccu_dataplane
		tlq_2718 天前改反了，统一用 asccomm_ccu.so
69	c_sec	69	c_sec
70	ascendcl	70	ascendcl
71	acl_rt	71	acl_rt

Mexamples/ccu/ccu_direct/01_allgather/main.asc+90-168

	@@ -20,16 +20,13 @@
20	#include <hccl/hccl_comm.h>	20	#include <hccl/hccl_comm.h>
21	#include <hccl/hccl_rank_graph.h>	21	#include <hccl/hccl_rank_graph.h>
22	#include <hccl/hccl_res.h>	22	#include <hccl/hccl_res.h>
23	-#include <hccl/hccl_types.h>	23	+#include <hcomm/hcomm_primitives.h>
24	-#include <hccl/hcomm_primitives.h>	24	+#include <ccu/hcomm/ccu_api_types.h>
25	-#include <ccu/ccu_types.h>	25	+#include <ccu/hcomm/ccu_resource_api.h>
26	-#include <ccu/ccu_variable.hpp>	26	+#include <ccu/hcomm/ccu_primitives.hpp>
27	-#include <ccu/ccu_event.hpp>
28	-#include <ccu/ccu_primitives.hpp>
29		27
30	-#include "hccl_ccu_res.h"	28	+#include <ccu/ccu_host_launch.h>
31	-#include "ccu_host_launch.h"	29	+#include "../../../common/ccu_instance.h"
32	-#include "ccu_res.h"
33		30
34	#ifdef ACLCHECK	31	#ifdef ACLCHECK
35	#undef ACLCHECK	32	#undef ACLCHECK
	@@ -89,14 +86,8 @@ constexpr uint32_t CCU_LOCAL_COPY_MS_PER_LOOP = 8;
89	constexpr uint32_t CCU_MS_LOCAL_COPY_LOOP_COUNT = 8;	86	constexpr uint32_t CCU_MS_LOCAL_COPY_LOOP_COUNT = 8;
90	constexpr uint64_t CCU_MAX_RANK_SIZE = 16;	87	constexpr uint64_t CCU_MAX_RANK_SIZE = 16;
91	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	88	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
92	-constexpr uint32_t HCOMM_HOST_TASK_ARG_NUM = 10;	89	+constexpr uint32_t CCU_DIRECT_TASK_ARG_NUM = 10;
		D D DragonBornHD8418 天前 [提示] 三个 ccu_direct 样例里同一个含义的常量两种叫法：01 改成了 CCU_DIRECT_TASK_ARG_NUM，02/03 还叫 HCOMM_HOST_TASK_ARG_NUM。<<<>>> 已经不走 hcomm 的 host task 链路了，建议 02/03 一起改掉，读样例不用来回对： `// 02 main.asc（值 10）、03 main.asc（值 8）里同步替换，删掉旧名 HCOMM_HOST_TASK_ARG_NUM constexpr uint32_t CCU_DIRECT_TASK_ARG_NUM = 10; // <- 02 样例；03 保持原值 8`
93	constexpr uint32_t CCU_DIE_NUM = 2;	90	constexpr uint32_t CCU_DIE_NUM = 2;
94	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_ADDRESS = 400;
95	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_LOOP = 16;
96	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_CCU_BUF = 128;
97	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_VARIABLE = 400;
98	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_EVENT = 48;
99	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_CCU_THREAD = 2;
100		91
101	constexpr uint32_t SIZE_TABLE[HCCL_DATA_TYPE_RESERVED] = {	92	constexpr uint32_t SIZE_TABLE[HCCL_DATA_TYPE_RESERVED] = {
102	sizeof(int8_t),	93	sizeof(int8_t),
	@@ -125,23 +116,23 @@ struct LoopGroupConfig {
125	};	116	};
126		117
127	struct LoopGroupResource {	118	struct LoopGroupResource {
128	- ccu::Array<ccu::Event> completedEvent{0};	119	+ ccu::array<ccu::event> completedEvent{0};
129	- ccu::Array<ccu::CcuBuffer> ccuBuf{0};	120	+ ccu::array<ccu::ccu_buffer> ccuBuf{0};
130	uint32_t eventCount;	121	uint32_t eventCount;
131	uint32_t bufCount;	122	uint32_t bufCount;
132	};	123	};
133		124
134	struct GroupOpSizeVars {	125	struct GroupOpSizeVars {
135	- ccu::Variable addrOffset;	126	+ ccu::variable addrOffset;
136	- ccu::Variable loopParam;	127	+ ccu::variable loopParam;
137	- ccu::Variable parallelParam;	128	+ ccu::variable parallelParam;
138	- ccu::Variable residual;	129	+ ccu::variable residual;
139	};	130	};
140		131
141	struct CcuLoopEntity {	132	struct CcuLoopEntity {
142	- std::unique_ptr<ccu::Func> body[2];	133	+ std::unique_ptr<ccu::func> body[2];
143	- std::unique_ptr<ccu::Loop> loops[2];	134	+ std::unique_ptr<ccu::loop> loops[2];
144	- ccu::Variable loopParam[2];	135	+ ccu::variable loopParam[2];
145	};	136	};
146		137
147	struct CcuKernelArgAllGatherMesh1DMem2Mem {	138	struct CcuKernelArgAllGatherMesh1DMem2Mem {
	@@ -154,13 +145,13 @@ struct CcuKernelArgAllGatherMesh1DMem2Mem {
154	struct AllGatherMesh1DMem2MemContext {	145	struct AllGatherMesh1DMem2MemContext {
155	const CcuKernelArgAllGatherMesh1DMem2Mem* arg;	146	const CcuKernelArgAllGatherMesh1DMem2Mem* arg;
156		147
157	- ccu::Variable input;	148	+ ccu::variable input;
158	- std::vector<ccu::Variable> output;	149	+ std::vector<ccu::variable> output;
159	- std::vector<ccu::Variable> token;	150	+ std::vector<ccu::variable> token;
160	- ccu::Variable currentRankSliceInputOffset;	151	+ ccu::variable currentRankSliceInputOffset;
161	- ccu::Variable currentRankSliceOutputOffset;	152	+ ccu::variable currentRankSliceOutputOffset;
162	- ccu::Variable sliceSize;	153	+ ccu::variable sliceSize;
163	- ccu::Event event;	154	+ ccu::event event;
164		155
165	GroupOpSizeVars goSize;	156	GroupOpSizeVars goSize;
166	LoopGroupConfig moConfig;	157	LoopGroupConfig moConfig;
	@@ -174,20 +165,6 @@ struct AllGatherMesh1DMem2MemContext {
174	bool IsLoopEntityRegistered(std::string loopStr) { return loopMap.count(loopStr) != 0; }	165	bool IsLoopEntityRegistered(std::string loopStr) { return loopMap.count(loopStr) != 0; }
175	};	166	};
176		167
177	-struct CcuDefaultRes {
178	- HcommCcuResType type;
179	- uint32_t num;
180	-};
181	-
182	-static constexpr CcuDefaultRes CCU_DEFAULT_RES[] = {
183	- {HCOMM_CCU_RES_TYPE_ADDRESS, CCU_DEFAULT_RES_FRACTION_ADDRESS},
184	- {HCOMM_CCU_RES_TYPE_LOOP, CCU_DEFAULT_RES_FRACTION_LOOP},
185	- {HCOMM_CCU_RES_TYPE_CCU_BUF, CCU_DEFAULT_RES_FRACTION_CCU_BUF},
186	- {HCOMM_CCU_RES_TYPE_VARIABLE, CCU_DEFAULT_RES_FRACTION_VARIABLE},
187	- {HCOMM_CCU_RES_TYPE_EVENT, CCU_DEFAULT_RES_FRACTION_EVENT},
188	- {HCOMM_CCU_RES_TYPE_CCU_THREAD, CCU_DEFAULT_RES_FRACTION_CCU_THREAD},
189	-};
190	-
191	static HcclResult ConvertCcuToHccl(CcuResult ccuResult)	168	static HcclResult ConvertCcuToHccl(CcuResult ccuResult)
192	{	169	{
193	switch (ccuResult) {	170	switch (ccuResult) {
	@@ -210,56 +187,6 @@ static HcclResult ConvertCcuToHccl(CcuResult ccuResult)
210	}	187	}
211	}	188	}
212		189
213	-static void DestroyAllDescs(std::vector<HcommCcuResDescHandle>& descs)
214	-{
215	- for (auto& desc : descs) {
216	- (void)HcommCcuInsResDescDestroy(desc);
217	- }
218	-}
219	-
220	-static HcclResult CreateCcuInsWithEnabledDies(CcuInsHandle& insHandle)
221	-{
222	- std::vector<HcommCcuResDescHandle> resDescs;
223	- for (uint32_t dieId = 0; dieId < CCU_DIE_NUM; dieId++) {
224	- HcommCcuResDescHandle desc = 0;
225	- CcuResult descCreateRet = HcommCcuInsResDescCreate(dieId, &desc);
226	- if (descCreateRet != CCU_SUCCESS) {
227	- (void)HcommCcuInsResDescDestroy(desc);
228	- DestroyAllDescs(resDescs);
229	- return ConvertCcuToHccl(descCreateRet);
230	- }
231	-
232	- CcuResult probeRet = HcommCcuQueryRemainResDesc(desc);
233	- if (probeRet == CCU_E_UNAVAIL) {
234	- (void)HcommCcuInsResDescDestroy(desc);
235	- continue;
236	- }
237	- if (probeRet != CCU_SUCCESS) {
238	- (void)HcommCcuInsResDescDestroy(desc);
239	- DestroyAllDescs(resDescs);
240	- return ConvertCcuToHccl(probeRet);
241	- }
242	-
243	- for (const auto& res : CCU_DEFAULT_RES) {
244	- CcuResult setRet = HcommCcuInsResDescSetNum(desc, res.type, res.num);
245	- if (setRet != CCU_SUCCESS) {
246	- (void)HcommCcuInsResDescDestroy(desc);
247	- DestroyAllDescs(resDescs);
248	- return ConvertCcuToHccl(setRet);
249	- }
250	- }
251	- resDescs.push_back(desc);
252	- }
253	-
254	- if (resDescs.empty()) {
255	- return HCCL_E_INTERNAL;
256	- }
257	-
258	- CcuResult createRet = HcommCcuInsCreate(resDescs.data(), static_cast<uint32_t>(resDescs.size()), &insHandle);
259	- DestroyAllDescs(resDescs);
260	- return ConvertCcuToHccl(createRet);
261	-}
262	-
263	static HcclResult AcquireCcuChannels(	190	static HcclResult AcquireCcuChannels(
264	HcclComm comm, uint32_t rankId, uint32_t rankSize, std::vector<ChannelHandle>& kernelChannels, uint32_t& dieId)	191	HcclComm comm, uint32_t rankId, uint32_t rankSize, std::vector<ChannelHandle>& kernelChannels, uint32_t& dieId)
265	{	192	{
	@@ -416,7 +343,7 @@ static uint64_t GetOffsetParam(uint64_t gsaOffset, uint64_t msOffset, uint64_t c
416	}	343	}
417		344
418	static void InitGroupCopyResources(	345	static void InitGroupCopyResources(
419	- AllGatherMesh1DMem2MemContext& ctx, ccu::LocalAddr* loopSrc, ccu::LocalAddr* loopDst, ccu::Variable* loopLen)	346	+ AllGatherMesh1DMem2MemContext& ctx, ccu::local_addr* loopSrc, ccu::local_addr* loopDst, ccu::variable* loopLen)
420	{	347	{
421	if (!ctx.resourceAllocated) {	348	if (!ctx.resourceAllocated) {
422	ctx.moConfig.msInterleave = CCU_MS_INTERLEAVE;	349	ctx.moConfig.msInterleave = CCU_MS_INTERLEAVE;
	@@ -424,10 +351,10 @@ static void InitGroupCopyResources(
424	ctx.moConfig.memSlice = CCU_LOCAL_COPY_MS_PER_LOOP * CCU_MS_SIZE;	351	ctx.moConfig.memSlice = CCU_LOCAL_COPY_MS_PER_LOOP * CCU_MS_SIZE;
425		352
426	ctx.moRes.eventCount = ctx.moConfig.loopCount;	353	ctx.moRes.eventCount = ctx.moConfig.loopCount;
427	- ctx.moRes.completedEvent = ccu::Array<ccu::Event>(ctx.moRes.eventCount);	354	+ ctx.moRes.completedEvent = ccu::array<ccu::event>(ctx.moRes.eventCount);
428		355
429	ctx.moRes.bufCount = ctx.moConfig.loopCount * ctx.moConfig.msInterleave;	356	ctx.moRes.bufCount = ctx.moConfig.loopCount * ctx.moConfig.msInterleave;
430	- ctx.moRes.ccuBuf = ccu::Array<ccu::CcuBuffer>(ctx.moRes.bufCount);	357	+ ctx.moRes.ccuBuf = ccu::array<ccu::ccu_buffer>(ctx.moRes.bufCount);
431		358
432	ctx.resourceAllocated = true;	359	ctx.resourceAllocated = true;
433	}	360	}
	@@ -438,24 +365,24 @@ static void InitGroupCopyResources(
438	auto& entity = ctx.loopMap[loopType];	365	auto& entity = ctx.loopMap[loopType];
439	for (uint32_t index = 0; index < 2; index++) {	366	for (uint32_t index = 0; index < 2; index++) {
440	uint32_t bufBase = index * ctx.moConfig.msInterleave;	367	uint32_t bufBase = index * ctx.moConfig.msInterleave;
441	- ccu::Event loopEvt = ctx.moRes.completedEvent[index];	368	+ ccu::event loopEvt = ctx.moRes.completedEvent[index];
442	- entity.body[index].reset(new ccu::Func([&ctx, index, bufBase, loopEvt, loopSrc, loopDst, loopLen]() {	369	+ entity.body[index].reset(new ccu::func([&ctx, index, bufBase, loopEvt, loopSrc, loopDst, loopLen]() {
443	ccu::LocalCopy(ctx.moRes.ccuBuf[bufBase], loopSrc[index], loopLen[index], loopEvt, 1);	370	ccu::LocalCopy(ctx.moRes.ccuBuf[bufBase], loopSrc[index], loopLen[index], loopEvt, 1);
444	- ccu::EventWait(loopEvt, 1);	371	+ ccu::event_wait(loopEvt, 1);
445	ccu::LocalCopy(loopDst[index], ctx.moRes.ccuBuf[bufBase], loopLen[index], loopEvt, 1);	372	ccu::LocalCopy(loopDst[index], ctx.moRes.ccuBuf[bufBase], loopLen[index], loopEvt, 1);
446	- ccu::EventWait(loopEvt, 1);	373	+ ccu::event_wait(loopEvt, 1);
447	}));	374	}));
448	- entity.loops[index].reset(new ccu::Loop(entity.loopParam[index], *entity.body[index]));	375	+ entity.loops[index].reset(new ccu::loop(entity.loopParam[index], *entity.body[index]));
449	}	376	}
450	}	377	}
451	}	378	}
452		379
453	static CcuResult GroupCopy(	380	static CcuResult GroupCopy(
454	- AllGatherMesh1DMem2MemContext& ctx, ccu::LocalAddr dst, ccu::LocalAddr src, GroupOpSizeVars goSize)	381	+ AllGatherMesh1DMem2MemContext& ctx, ccu::local_addr dst, ccu::local_addr src, GroupOpSizeVars goSize)
455	{	382	{
456	- ccu::LocalAddr loopSrc[2];	383	+ ccu::local_addr loopSrc[2];
457	- ccu::LocalAddr loopDst[2];	384	+ ccu::local_addr loopDst[2];
458	- ccu::Variable loopLen[2];	385	+ ccu::variable loopLen[2];
459		386
460	InitGroupCopyResources(ctx, loopSrc, loopDst, loopLen);	387	InitGroupCopyResources(ctx, loopSrc, loopDst, loopLen);
461		388
	@@ -463,63 +390,63 @@ static CcuResult GroupCopy(
463		390
464	CCU_IF(goSize.addrOffset != 0)	391	CCU_IF(goSize.addrOffset != 0)
465	{	392	{
466	- ccu::Variable loopParam;	393	+ ccu::variable loopParam;
467	loopParam = GetLoopParam(0, ctx.moConfig.memSlice * ctx.moConfig.loopCount, 0);	394	loopParam = GetLoopParam(0, ctx.moConfig.memSlice * ctx.moConfig.loopCount, 0);
468	loopParam += goSize.loopParam;	395	loopParam += goSize.loopParam;
469		396
470	- ccu::Variable sliceSize;	397	+ ccu::variable sliceSize;
471	sliceSize = ctx.moConfig.memSlice;	398	sliceSize = ctx.moConfig.memSlice;
472		399
473	- loopSrc[0].addr = src.addr;	400	+ loopSrc[0].addr_ = src.addr_;
474	loopSrc[0].token = src.token;	401	loopSrc[0].token = src.token;
475	- loopDst[0].addr = dst.addr;	402	+ loopDst[0].addr_ = dst.addr_;
476	loopDst[0].token = dst.token;	403	loopDst[0].token = dst.token;
477	loopLen[0] = sliceSize;	404	loopLen[0] = sliceSize;
478		405
479	loops.loopParam[0] = loopParam;	406	loops.loopParam[0] = loopParam;
480	- ccu::Variable paraCfg;	407	+ ccu::variable paraCfg;
481	paraCfg = GetParallelParam(ctx.moConfig.loopCount - 1, 0, 1);	408	paraCfg = GetParallelParam(ctx.moConfig.loopCount - 1, 0, 1);
482		409
483	- ccu::Variable offsetCfg;	410	+ ccu::variable offsetCfg;
484	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);	411	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);
485	- std::vector<ccu::Loop> grpLoops{*loops.loops[0]};	412	+ std::vector<ccu::loop> grpLoops{*loops.loops[0]};
486	- ccu::LoopGroup group(paraCfg, offsetCfg, ctx.moConfig.loopCount, grpLoops);	413	+ ccu::loop_group group(paraCfg, offsetCfg, ctx.moConfig.loopCount, grpLoops);
487	}	414	}
488		415
489	CCU_IF(goSize.parallelParam != 0)	416	CCU_IF(goSize.parallelParam != 0)
490	{	417	{
491	- src.addr += goSize.addrOffset;	418	+ src.addr_ += goSize.addrOffset;
492	- dst.addr += goSize.addrOffset;	419	+ dst.addr_ += goSize.addrOffset;
493		420
494	- loopSrc[0].addr = src.addr;	421	+ loopSrc[0].addr_ = src.addr_;
495	loopSrc[0].token = src.token;	422	loopSrc[0].token = src.token;
496	- loopDst[0].addr = dst.addr;	423	+ loopDst[0].addr_ = dst.addr_;
497	loopDst[0].token = dst.token;	424	loopDst[0].token = dst.token;
498	loopLen[0] = goSize.residual;	425	loopLen[0] = goSize.residual;
499		426
500	- src.addr += goSize.residual;	427	+ src.addr_ += goSize.residual;
501	- dst.addr += goSize.residual;	428	+ dst.addr_ += goSize.residual;
502		429
503	- ccu::Variable sliceSize;	430	+ ccu::variable sliceSize;
504	sliceSize = ctx.moConfig.memSlice;	431	sliceSize = ctx.moConfig.memSlice;
505		432
506	- loopSrc[1].addr = src.addr;	433	+ loopSrc[1].addr_ = src.addr_;
507	loopSrc[1].token = src.token;	434	loopSrc[1].token = src.token;
508	- loopDst[1].addr = dst.addr;	435	+ loopDst[1].addr_ = dst.addr_;
509	loopDst[1].token = dst.token;	436	loopDst[1].token = dst.token;
510	loopLen[1] = sliceSize;	437	loopLen[1] = sliceSize;
511		438
512	- ccu::Variable loopCfg0;	439	+ ccu::variable loopCfg0;
513	loopCfg0 = GetLoopParam(0, 0, 1);	440	loopCfg0 = GetLoopParam(0, 0, 1);
514	- ccu::Variable loopCfg1;	441	+ ccu::variable loopCfg1;
515	loopCfg1 = GetLoopParam(0, 0, 1);	442	loopCfg1 = GetLoopParam(0, 0, 1);
516	- ccu::Variable offsetCfg;	443	+ ccu::variable offsetCfg;
517	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);	444	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);
518		445
519	loops.loopParam[0] = loopCfg0;	446	loops.loopParam[0] = loopCfg0;
520	loops.loopParam[1] = loopCfg1;	447	loops.loopParam[1] = loopCfg1;
521	- std::vector<ccu::Loop> grpLoops{loops.loops[0], loops.loops[1]};	448	+ std::vector<ccu::loop> grpLoops{loops.loops[0], loops.loops[1]};
522	- ccu::LoopGroup group(goSize.parallelParam, offsetCfg, ctx.moConfig.loopCount, grpLoops);	449	+ ccu::loop_group group(goSize.parallelParam, offsetCfg, ctx.moConfig.loopCount, grpLoops);
523	}	450	}
524		451
525	return CCU_SUCCESS;	452	return CCU_SUCCESS;
	@@ -527,24 +454,24 @@ static CcuResult GroupCopy(
527		454
528	static CcuResult ExecuteAllGatherTransfer(AllGatherMesh1DMem2MemContext& ctx)	455	static CcuResult ExecuteAllGatherTransfer(AllGatherMesh1DMem2MemContext& ctx)
529	{	456	{
530	- ccu::LocalAddr src;	457	+ ccu::local_addr src;
531	- ccu::LocalAddr localDst;	458	+ ccu::local_addr localDst;
532	- std::vector<ccu::RemoteAddr> dst;	459	+ std::vector<ccu::remote_addr> dst;
533		460
534	dst.resize(ctx.arg->rankSize);	461	dst.resize(ctx.arg->rankSize);
535		462
536	- src.addr = ctx.input;	463	+ src.addr_ = ctx.input;
537	- src.addr += ctx.currentRankSliceInputOffset;	464	+ src.addr_ += ctx.currentRankSliceInputOffset;
538	src.token = ctx.token[ctx.arg->rankId];	465	src.token = ctx.token[ctx.arg->rankId];
539		466
540	for (uint32_t rankIdx = 0; rankIdx < ctx.arg->rankSize; rankIdx++) {	467	for (uint32_t rankIdx = 0; rankIdx < ctx.arg->rankSize; rankIdx++) {
541	if (rankIdx == ctx.arg->rankId) {	468	if (rankIdx == ctx.arg->rankId) {
542	- localDst.addr = ctx.output[rankIdx];	469	+ localDst.addr_ = ctx.output[rankIdx];
543	- localDst.addr += ctx.currentRankSliceOutputOffset;	470	+ localDst.addr_ += ctx.currentRankSliceOutputOffset;
544	localDst.token = ctx.token[rankIdx];	471	localDst.token = ctx.token[rankIdx];
545	} else {	472	} else {
546	- dst[rankIdx].addr = ctx.output[rankIdx];	473	+ dst[rankIdx].addr_ = ctx.output[rankIdx];
547	- dst[rankIdx].addr += ctx.currentRankSliceOutputOffset;	474	+ dst[rankIdx].addr_ += ctx.currentRankSliceOutputOffset;
548	dst[rankIdx].token = ctx.token[rankIdx];	475	dst[rankIdx].token = ctx.token[rankIdx];
549	}	476	}
550	}	477	}
	@@ -563,17 +490,17 @@ static CcuResult ExecuteAllGatherTransfer(AllGatherMesh1DMem2MemContext& ctx)
563	}	490	}
564		491
565	CCU_CHK_RET(GroupCopy(ctx, localDst, src, ctx.goSize));	492	CCU_CHK_RET(GroupCopy(ctx, localDst, src, ctx.goSize));
566	- CCU_CHK_RET(ccu::EventRecord(ctx.event, 1 << ctx.arg->rankId));	493	+ CCU_CHK_RET(ccu::event_record(ctx.event, 1 << ctx.arg->rankId));
567		494
568	const uint16_t totalMask = (1 << ctx.arg->rankSize) - 1;	495	const uint16_t totalMask = (1 << ctx.arg->rankSize) - 1;
569	- CCU_CHK_RET(ccu::EventWait(ctx.event, totalMask));	496	+ CCU_CHK_RET(ccu::event_wait(ctx.event, totalMask));
570		497
571	return CcuResult::CCU_SUCCESS;	498	return CcuResult::CCU_SUCCESS;
572	}	499	}
573		500
574	__ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(	501	__ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(
575	uint64_t sqeArgs1, uint64_t sqeArgs2, uint64_t sqeArgs3, uint64_t sqeArgs4, uint64_t sqeArgs5, uint64_t sqeArgs6,	502	uint64_t sqeArgs1, uint64_t sqeArgs2, uint64_t sqeArgs3, uint64_t sqeArgs4, uint64_t sqeArgs5, uint64_t sqeArgs6,
576	- uint64_t sqeArgs7, uint64_t sqeArgs8, uint64_t sqeArgs9, uint64_t sqeArgs10, CcuKernelArg arg)	503	+ uint64_t sqeArgs7, uint64_t sqeArgs8, uint64_t sqeArgs9, uint64_t sqeArgs10, ccu_kernel_arg arg)
577	{	504	{
578	auto* kernelArg = static_cast<CcuKernelArgAllGatherMesh1DMem2Mem*>(arg);	505	auto* kernelArg = static_cast<CcuKernelArgAllGatherMesh1DMem2Mem*>(arg);
579		506
	@@ -590,43 +517,43 @@ __ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(
590	uint32_t channelIdx = 0;	517	uint32_t channelIdx = 0;
591	for (uint64_t peerId = 0; peerId < ctx.arg->rankSize; peerId++) {	518	for (uint64_t peerId = 0; peerId < ctx.arg->rankSize; peerId++) {
592	if (peerId != ctx.arg->rankId) {	519	if (peerId != ctx.arg->rankId) {
593	- ctx.output[peerId] = ccu::GetResByChannel<ccu::Variable>(ctx.arg->channels[channelIdx], OUTPUT_XN_ID);	520	+ ctx.output[peerId] = ccu::GetResByChannel<ccu::variable>(ctx.arg->channels[channelIdx], OUTPUT_XN_ID);
594	- ctx.token[peerId] = ccu::GetResByChannel<ccu::Variable>(ctx.arg->channels[channelIdx], TOKEN_XN_ID);	521	+ ctx.token[peerId] = ccu::GetResByChannel<ccu::variable>(ctx.arg->channels[channelIdx], TOKEN_XN_ID);
595	channelIdx++;	522	channelIdx++;
596	}	523	}
597	}	524	}
598		525
599	uint32_t argId = 0;	526	uint32_t argId = 0;
600	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.input, argId++));	527	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.input, argId++));
601	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.output[ctx.arg->rankId], argId++));	528	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.output[ctx.arg->rankId], argId++));
602	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.token[ctx.arg->rankId], argId++));	529	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.token[ctx.arg->rankId], argId++));
603	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.currentRankSliceInputOffset, argId++));	530	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.currentRankSliceInputOffset, argId++));
604	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.currentRankSliceOutputOffset, argId++));	531	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.currentRankSliceOutputOffset, argId++));
605	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.sliceSize, argId++));	532	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.sliceSize, argId++));
606	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.goSize.addrOffset, argId++));	533	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.goSize.addrOffset, argId++));
607	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.goSize.loopParam, argId++));	534	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.goSize.loopParam, argId++));
608	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.goSize.parallelParam, argId++));	535	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.goSize.parallelParam, argId++));
609	- CCU_CHK_VOID_RET(ccu::LoadArg(ctx.goSize.residual, argId++));	536	+ CCU_CHK_VOID_RET(ccu::load_arg(ctx.goSize.residual, argId++));
610		537
611	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	538	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
612	- CCU_CHK_VOID_RET(ccu::WriteVariableWithNotify(	539	+ CCU_CHK_VOID_RET(ccu::write_variable_with_notify(
613	ctx.arg->channels[i], ctx.output[ctx.arg->rankId], OUTPUT_XN_ID, CKE_IDX_0, 1 << OUTPUT_XN_ID));	540	ctx.arg->channels[i], ctx.output[ctx.arg->rankId], OUTPUT_XN_ID, CKE_IDX_0, 1 << OUTPUT_XN_ID));
614	- CCU_CHK_VOID_RET(ccu::WriteVariableWithNotify(	541	+ CCU_CHK_VOID_RET(ccu::write_variable_with_notify(
615	ctx.arg->channels[i], ctx.token[ctx.arg->rankId], TOKEN_XN_ID, CKE_IDX_0, 1 << TOKEN_XN_ID));	542	ctx.arg->channels[i], ctx.token[ctx.arg->rankId], TOKEN_XN_ID, CKE_IDX_0, 1 << TOKEN_XN_ID));
616	}	543	}
617		544
618	uint32_t allBit = (1 << OUTPUT_XN_ID) \| (1 << TOKEN_XN_ID);	545	uint32_t allBit = (1 << OUTPUT_XN_ID) \| (1 << TOKEN_XN_ID);
619	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	546	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
620	- CCU_CHK_VOID_RET(ccu::NotifyWait(ctx.arg->channels[i], CKE_IDX_0, allBit));	547	+ CCU_CHK_VOID_RET(ccu::notify_wait(ctx.arg->channels[i], CKE_IDX_0, allBit));
621	}	548	}
622		549
623	CCU_CHK_VOID_RET(ExecuteAllGatherTransfer(ctx));	550	CCU_CHK_VOID_RET(ExecuteAllGatherTransfer(ctx));
624		551
625	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	552	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
626	- CCU_CHK_VOID_RET(ccu::NotifyRecord(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));	553	+ CCU_CHK_VOID_RET(ccu::notify_record(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));
627	}	554	}
628	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	555	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
629	- CCU_CHK_VOID_RET(ccu::NotifyWait(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));	556	+ CCU_CHK_VOID_RET(ccu::notify_wait(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));
630	}	557	}
631	}	558	}
632		559
	@@ -670,18 +597,13 @@ HcclResult AllGatherCustom(
670	}	597	}
671	const uint64_t dieMask = 1ULL << dieId;	598	const uint64_t dieMask = 1ULL << dieId;
672		599
673	- CcuInsHandle insHandle{0};	600	+ CcuInsHandle insHandle{};
674	- HCCLCHECK(CreateCcuInsWithEnabledDies(insHandle));	601	+ HCCLCHECK(asccomm_examples::QueryCcuInstance(comm, insHandle));
675	- HcclResult assignRet = HcclCommAssignCcuIns(comm, insHandle);
676	- if (assignRet != HCCL_SUCCESS) {
677	- (void)HcommCcuInsDestroy(insHandle);
678	- return assignRet;
679	- }
680		602
681	uint64_t token = 0;	603	uint64_t token = 0;
682	uint64_t baseInputAddr = reinterpret_cast<uint64_t>(sendBuf);	604	uint64_t baseInputAddr = reinterpret_cast<uint64_t>(sendBuf);
683	uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(recvBuf);	605	uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(recvBuf);
684	- HCCLCHECK(HcommCcuGetMemToken(baseInputAddr, inputSize, &token));	606	+ HCCLCHECK(ConvertCcuToHccl(asccomm_ccu_get_mem_token(baseInputAddr, inputSize, &token)));
685		607
686	LoopGroupConfig config{};	608	LoopGroupConfig config{};
687	config.msInterleave = CCU_MS_INTERLEAVE;	609	config.msInterleave = CCU_MS_INTERLEAVE;
	@@ -689,7 +611,7 @@ HcclResult AllGatherCustom(
689	config.memSlice = CCU_MS_SIZE * CCU_LOCAL_COPY_MS_PER_LOOP;	611	config.memSlice = CCU_MS_SIZE * CCU_LOCAL_COPY_MS_PER_LOOP;
690	auto goSize = CalGoSize(inputSize, config);	612	auto goSize = CalGoSize(inputSize, config);
691		613
692	- uint64_t taskArgs[HCOMM_HOST_TASK_ARG_NUM] = {	614	+ uint64_t taskArgs[CCU_DIRECT_TASK_ARG_NUM] = {
693	baseInputAddr, baseOutputAddr, token, 0, inputSize * rankId,	615	baseInputAddr, baseOutputAddr, token, 0, inputSize * rankId,
694	inputSize, goSize[0], goSize[1], goSize[2], goSize[3],	616	inputSize, goSize[0], goSize[1], goSize[2], goSize[3],
695	};	617	};

Mexamples/ccu/ccu_direct/02_allgather_add/CMakeLists.txt+5-5

	@@ -30,12 +30,10 @@ target_include_directories(demo PRIVATE
30	# CANN Toolkit头文件	30	# CANN Toolkit头文件
31	${ASCEND_CANN_PACKAGE_PATH}/include	31	${ASCEND_CANN_PACKAGE_PATH}/include
32	${ASCEND_CANN_PACKAGE_PATH}/include/hccl	32	${ASCEND_CANN_PACKAGE_PATH}/include/hccl
33	- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm	33	+ ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api
34	- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm/ccu
35	- ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api/ccu
36	${ASCEND_CANN_PACKAGE_PATH}/pkg_inc	34	${ASCEND_CANN_PACKAGE_PATH}/pkg_inc
37	- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/hcomm	35	+ # 样例公共辅助头（ccu_instance.h）
38	- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/hcomm/ccu	36	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../common
39	)	37	)
40		38
41	target_compile_features(demo PRIVATE cxx_std_14)	39	target_compile_features(demo PRIVATE cxx_std_14)
	@@ -59,6 +57,7 @@ target_link_options(demo PRIVATE
59	-Wl,-z,relro	57	-Wl,-z,relro
60	-Wl,-z,now	58	-Wl,-z,now
61	-Wl,-z,noexecstack	59	-Wl,-z,noexecstack
		60	+ -Wl,-rpath,${ASCEND_CANN_PACKAGE_PATH}/lib64
62	)	61	)
63	target_link_directories(demo PRIVATE	62	target_link_directories(demo PRIVATE
64	${ASCEND_CANN_PACKAGE_PATH}/lib64	63	${ASCEND_CANN_PACKAGE_PATH}/lib64
	@@ -66,6 +65,7 @@ target_link_directories(demo PRIVATE
66	target_link_libraries(demo PRIVATE	65	target_link_libraries(demo PRIVATE
67	hcomm	66	hcomm
68	asccomm_ccu	67	asccomm_ccu
		68	+ asccomm_ccu_dataplane
69	c_sec	69	c_sec
70	ascendcl	70	ascendcl
71	acl_rt	71	acl_rt

Mexamples/ccu/ccu_direct/02_allgather_add/main.asc+100-179

	@@ -20,16 +20,13 @@
20	#include <hccl/hccl_comm.h>	20	#include <hccl/hccl_comm.h>
21	#include <hccl/hccl_rank_graph.h>	21	#include <hccl/hccl_rank_graph.h>
22	#include <hccl/hccl_res.h>	22	#include <hccl/hccl_res.h>
23	-#include <hccl/hccl_types.h>	23	+#include <hcomm/hcomm_primitives.h>
24	-#include <hccl/hcomm_primitives.h>	24	+#include <ccu/hcomm/ccu_api_types.h>
25	-#include <ccu/ccu_types.h>	25	+#include <ccu/hcomm/ccu_resource_api.h>
26	-#include <ccu/ccu_variable.hpp>	26	+#include <ccu/hcomm/ccu_primitives.hpp>
27	-#include <ccu/ccu_event.hpp>
28	-#include <ccu/ccu_primitives.hpp>
29		27
30	-#include "hccl_ccu_res.h"	28	+#include <ccu/ccu_host_launch.h>
31	-#include "ccu_host_launch.h"	29	+#include "../../../common/ccu_instance.h"
32	-#include "ccu_res.h"
33	#include "kernel_operator.h"	30	#include "kernel_operator.h"
34		31
35	#ifdef ACLCHECK	32	#ifdef ACLCHECK
	@@ -90,15 +87,9 @@ constexpr uint32_t CCU_LOCAL_COPY_MS_PER_LOOP = 8;
90	constexpr uint32_t CCU_MS_LOCAL_COPY_LOOP_COUNT = 8;	87	constexpr uint32_t CCU_MS_LOCAL_COPY_LOOP_COUNT = 8;
91	constexpr uint64_t CCU_MAX_RANK_SIZE = 16;	88	constexpr uint64_t CCU_MAX_RANK_SIZE = 16;
92	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	89	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
93	-constexpr uint32_t HCOMM_HOST_TASK_ARG_NUM = 10;	90	+constexpr uint32_t CCU_DIRECT_TASK_ARG_NUM = 10;
94	constexpr uint32_t CKE_NUM = 2;	91	constexpr uint32_t CKE_NUM = 2;
95	constexpr uint32_t CCU_DIE_NUM = 2;	92	constexpr uint32_t CCU_DIE_NUM = 2;
96	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_ADDRESS = 400;
97	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_LOOP = 16;
98	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_CCU_BUF = 128;
99	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_VARIABLE = 400;
100	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_EVENT = 48;
101	-constexpr uint32_t CCU_DEFAULT_RES_FRACTION_CCU_THREAD = 2;
102		93
103	constexpr uint32_t SIZE_TABLE[HCCL_DATA_TYPE_RESERVED] = {	94	constexpr uint32_t SIZE_TABLE[HCCL_DATA_TYPE_RESERVED] = {
104	sizeof(int8_t),	95	sizeof(int8_t),
	@@ -127,23 +118,23 @@ struct LoopGroupConfig {
127	};	118	};
128		119
129	struct LoopGroupResource {	120	struct LoopGroupResource {
130	- ccu::Array<ccu::Event> completedEvent{0};	121	+ ccu::array<ccu::event> completedEvent{0};
131	- ccu::Array<ccu::CcuBuffer> ccuBuf{0};	122	+ ccu::array<ccu::ccu_buffer> ccuBuf{0};
132	uint32_t eventCount;	123	uint32_t eventCount;
133	uint32_t bufCount;	124	uint32_t bufCount;
134	};	125	};
135		126
136	struct GroupOpSizeVars {	127	struct GroupOpSizeVars {
137	- ccu::Variable addrOffset;	128	+ ccu::variable addrOffset;
138	- ccu::Variable loopParam;	129	+ ccu::variable loopParam;
139	- ccu::Variable parallelParam;	130	+ ccu::variable parallelParam;
140	- ccu::Variable residual;	131	+ ccu::variable residual;
141	};	132	};
142		133
143	struct CcuLoopEntity {	134	struct CcuLoopEntity {
144	- std::unique_ptr<ccu::Func> body[2];	135	+ std::unique_ptr<ccu::func> body[2];
145	- std::unique_ptr<ccu::Loop> loops[2];	136	+ std::unique_ptr<ccu::loop> loops[2];
146	- ccu::Variable loopParam[2];	137	+ ccu::variable loopParam[2];
147	};	138	};
148		139
149	struct CcuKernelArgAllGatherMesh1DMem2Mem {	140	struct CcuKernelArgAllGatherMesh1DMem2Mem {
	@@ -151,22 +142,27 @@ struct CcuKernelArgAllGatherMesh1DMem2Mem {
151	uint32_t channelCount;	142	uint32_t channelCount;
152	uint64_t rankSize;	143	uint64_t rankSize;
153	uint32_t rankId;	144	uint32_t rankId;
154	- CcuVariableHandle varHandle;	145	+ ccu_variable_handle varHandle;
155	uint32_t varNum;	146	uint32_t varNum;
156	- CcuEventHandle eventHandle;	147	+ ccu_event_handle eventHandle;
157	uint32_t eventNum;	148	uint32_t eventNum;
158	};	149	};
159		150
		151	+struct CcuMappedAddrs {
		152	+ uint64_t eventVa[CKE_NUM];
		153	+ uint64_t startVa[CCU_DIRECT_TASK_ARG_NUM];
		154	+};
		155	+
160	struct AllGatherMesh1DMem2MemContext {	156	struct AllGatherMesh1DMem2MemContext {
161	const CcuKernelArgAllGatherMesh1DMem2Mem* arg;	157	const CcuKernelArgAllGatherMesh1DMem2Mem* arg;
162		158
163	- ccu::Variable input;	159	+ ccu::variable input;
164	- std::vector<ccu::Variable> output;	160	+ std::vector<ccu::variable> output;
165	- std::vector<ccu::Variable> token;	161	+ std::vector<ccu::variable> token;
166	- ccu::Variable currentRankSliceInputOffset;	162	+ ccu::variable currentRankSliceInputOffset;
167	- ccu::Variable currentRankSliceOutputOffset;	163	+ ccu::variable currentRankSliceOutputOffset;
168	- ccu::Variable sliceSize;	164	+ ccu::variable sliceSize;
169	- ccu::Event event;	165	+ ccu::event event;
170		166
171	GroupOpSizeVars goSize;	167	GroupOpSizeVars goSize;
172	LoopGroupConfig moConfig;	168	LoopGroupConfig moConfig;
	@@ -180,25 +176,6 @@ struct AllGatherMesh1DMem2MemContext {
180	bool IsLoopEntityRegistered(std::string loopStr) { return loopMap.count(loopStr) != 0; }	176	bool IsLoopEntityRegistered(std::string loopStr) { return loopMap.count(loopStr) != 0; }
181	};	177	};
182		178
183	-struct CcuDefaultRes {
184	- HcommCcuResType type;
185	- uint32_t num;
186	-};
187	-
188	-struct CcuMappedAddrs {
189	- uint64_t eventVa[CKE_NUM];
190	- uint64_t startVa[HCOMM_HOST_TASK_ARG_NUM];
191	-};
192	-
193	-static constexpr CcuDefaultRes CCU_DEFAULT_RES[] = {
194	- {HCOMM_CCU_RES_TYPE_ADDRESS, CCU_DEFAULT_RES_FRACTION_ADDRESS},
195	- {HCOMM_CCU_RES_TYPE_LOOP, CCU_DEFAULT_RES_FRACTION_LOOP},
196	- {HCOMM_CCU_RES_TYPE_CCU_BUF, CCU_DEFAULT_RES_FRACTION_CCU_BUF},
197	- {HCOMM_CCU_RES_TYPE_VARIABLE, CCU_DEFAULT_RES_FRACTION_VARIABLE},
198	- {HCOMM_CCU_RES_TYPE_EVENT, CCU_DEFAULT_RES_FRACTION_EVENT},
199	- {HCOMM_CCU_RES_TYPE_CCU_THREAD, CCU_DEFAULT_RES_FRACTION_CCU_THREAD},
200	-};
201	-
202	static HcclResult ConvertCcuToHccl(CcuResult ccuResult)	179	static HcclResult ConvertCcuToHccl(CcuResult ccuResult)
203	{	180	{
204	switch (ccuResult) {	181	switch (ccuResult) {
	@@ -221,56 +198,6 @@ static HcclResult ConvertCcuToHccl(CcuResult ccuResult)
221	}	198	}
222	}	199	}
223		200
224	-static void DestroyAllDescs(std::vector<HcommCcuResDescHandle>& descs)
225	-{
226	- for (auto& desc : descs) {
227	- (void)HcommCcuInsResDescDestroy(desc);
228	- }
229	-}
230	-
231	-static HcclResult CreateCcuInsWithEnabledDies(CcuInsHandle& insHandle)
232	-{
233	- std::vector<HcommCcuResDescHandle> resDescs;
234	- for (uint32_t dieId = 0; dieId < CCU_DIE_NUM; dieId++) {
235	- HcommCcuResDescHandle desc = 0;
236	- CcuResult descCreateRet = HcommCcuInsResDescCreate(dieId, &desc);
237	- if (descCreateRet != CCU_SUCCESS) {
238	- (void)HcommCcuInsResDescDestroy(desc);
239	- DestroyAllDescs(resDescs);
240	- return ConvertCcuToHccl(descCreateRet);
241	- }
242	-
243	- CcuResult probeRet = HcommCcuQueryRemainResDesc(desc);
244	- if (probeRet == CCU_E_UNAVAIL) {
245	- (void)HcommCcuInsResDescDestroy(desc);
246	- continue;
247	- }
248	- if (probeRet != CCU_SUCCESS) {
249	- (void)HcommCcuInsResDescDestroy(desc);
250	- DestroyAllDescs(resDescs);
251	- return ConvertCcuToHccl(probeRet);
252	- }
253	-
254	- for (const auto& res : CCU_DEFAULT_RES) {
255	- CcuResult setRet = HcommCcuInsResDescSetNum(desc, res.type, res.num);
256	- if (setRet != CCU_SUCCESS) {
257	- (void)HcommCcuInsResDescDestroy(desc);
258	- DestroyAllDescs(resDescs);
259	- return ConvertCcuToHccl(setRet);
260	- }
261	- }
262	- resDescs.push_back(desc);
263	- }
264	-
265	- if (resDescs.empty()) {
266	- return HCCL_E_INTERNAL;
267	- }
268	-
269	- CcuResult createRet = HcommCcuInsCreate(resDescs.data(), static_cast<uint32_t>(resDescs.size()), &insHandle);
270	- DestroyAllDescs(resDescs);
271	- return ConvertCcuToHccl(createRet);
272	-}
273	-
274	static HcclResult AcquireCcuChannels(	201	static HcclResult AcquireCcuChannels(
275	HcclComm comm, uint32_t rankId, uint32_t rankSize, std::vector<ChannelHandle>& kernelChannels, uint32_t& dieId)	202	HcclComm comm, uint32_t rankId, uint32_t rankSize, std::vector<ChannelHandle>& kernelChannels, uint32_t& dieId)
276	{	203	{
	@@ -427,7 +354,7 @@ static uint64_t GetOffsetParam(uint64_t gsaOffset, uint64_t msOffset, uint64_t c
427	}	354	}
428		355
429	static void InitGroupCopyResources(	356	static void InitGroupCopyResources(
430	- AllGatherMesh1DMem2MemContext& ctx, ccu::LocalAddr* loopSrc, ccu::LocalAddr* loopDst, ccu::Variable* loopLen)	357	+ AllGatherMesh1DMem2MemContext& ctx, ccu::local_addr* loopSrc, ccu::local_addr* loopDst, ccu::variable* loopLen)
431	{	358	{
432	if (!ctx.resourceAllocated) {	359	if (!ctx.resourceAllocated) {
433	ctx.moConfig.msInterleave = CCU_MS_INTERLEAVE;	360	ctx.moConfig.msInterleave = CCU_MS_INTERLEAVE;
	@@ -435,10 +362,10 @@ static void InitGroupCopyResources(
435	ctx.moConfig.memSlice = CCU_LOCAL_COPY_MS_PER_LOOP * CCU_MS_SIZE;	362	ctx.moConfig.memSlice = CCU_LOCAL_COPY_MS_PER_LOOP * CCU_MS_SIZE;
436		363
437	ctx.moRes.eventCount = ctx.moConfig.loopCount;	364	ctx.moRes.eventCount = ctx.moConfig.loopCount;
438	- ctx.moRes.completedEvent = ccu::Array<ccu::Event>(ctx.moRes.eventCount);	365	+ ctx.moRes.completedEvent = ccu::array<ccu::event>(ctx.moRes.eventCount);
439		366
440	ctx.moRes.bufCount = ctx.moConfig.loopCount * ctx.moConfig.msInterleave;	367	ctx.moRes.bufCount = ctx.moConfig.loopCount * ctx.moConfig.msInterleave;
441	- ctx.moRes.ccuBuf = ccu::Array<ccu::CcuBuffer>(ctx.moRes.bufCount);	368	+ ctx.moRes.ccuBuf = ccu::array<ccu::ccu_buffer>(ctx.moRes.bufCount);
442		369
443	ctx.resourceAllocated = true;	370	ctx.resourceAllocated = true;
444	}	371	}
	@@ -449,24 +376,24 @@ static void InitGroupCopyResources(
449	auto& entity = ctx.loopMap[loopType];	376	auto& entity = ctx.loopMap[loopType];
450	for (uint32_t index = 0; index < 2; index++) {	377	for (uint32_t index = 0; index < 2; index++) {
451	uint32_t bufBase = index * ctx.moConfig.msInterleave;	378	uint32_t bufBase = index * ctx.moConfig.msInterleave;
452	- ccu::Event loopEvt = ctx.moRes.completedEvent[index];	379	+ ccu::event loopEvt = ctx.moRes.completedEvent[index];
453	- entity.body[index].reset(new ccu::Func([&ctx, index, bufBase, loopEvt, loopSrc, loopDst, loopLen]() {	380	+ entity.body[index].reset(new ccu::func([&ctx, index, bufBase, loopEvt, loopSrc, loopDst, loopLen]() {
454	ccu::LocalCopy(ctx.moRes.ccuBuf[bufBase], loopSrc[index], loopLen[index], loopEvt, 1);	381	ccu::LocalCopy(ctx.moRes.ccuBuf[bufBase], loopSrc[index], loopLen[index], loopEvt, 1);
455	- ccu::EventWait(loopEvt, 1);	382	+ ccu::event_wait(loopEvt, 1);
456	ccu::LocalCopy(loopDst[index], ctx.moRes.ccuBuf[bufBase], loopLen[index], loopEvt, 1);	383	ccu::LocalCopy(loopDst[index], ctx.moRes.ccuBuf[bufBase], loopLen[index], loopEvt, 1);
457	- ccu::EventWait(loopEvt, 1);	384	+ ccu::event_wait(loopEvt, 1);
458	}));	385	}));
459	- entity.loops[index].reset(new ccu::Loop(entity.loopParam[index], *entity.body[index]));	386	+ entity.loops[index].reset(new ccu::loop(entity.loopParam[index], *entity.body[index]));
460	}	387	}
461	}	388	}
462	}	389	}
463		390
464	static CcuResult GroupCopy(	391	static CcuResult GroupCopy(
465	- AllGatherMesh1DMem2MemContext& ctx, ccu::LocalAddr dst, ccu::LocalAddr src, GroupOpSizeVars goSize)	392	+ AllGatherMesh1DMem2MemContext& ctx, ccu::local_addr dst, ccu::local_addr src, GroupOpSizeVars goSize)
466	{	393	{
467	- ccu::LocalAddr loopSrc[2];	394	+ ccu::local_addr loopSrc[2];
468	- ccu::LocalAddr loopDst[2];	395	+ ccu::local_addr loopDst[2];
469	- ccu::Variable loopLen[2];	396	+ ccu::variable loopLen[2];
470		397
471	InitGroupCopyResources(ctx, loopSrc, loopDst, loopLen);	398	InitGroupCopyResources(ctx, loopSrc, loopDst, loopLen);
472		399
	@@ -474,63 +401,63 @@ static CcuResult GroupCopy(
474		401
475	CCU_IF(goSize.addrOffset != 0)	402	CCU_IF(goSize.addrOffset != 0)
476	{	403	{
477	- ccu::Variable loopParam;	404	+ ccu::variable loopParam;
478	loopParam = GetLoopParam(0, ctx.moConfig.memSlice * ctx.moConfig.loopCount, 0);	405	loopParam = GetLoopParam(0, ctx.moConfig.memSlice * ctx.moConfig.loopCount, 0);
479	loopParam += goSize.loopParam;	406	loopParam += goSize.loopParam;
480		407
481	- ccu::Variable sliceSize;	408	+ ccu::variable sliceSize;
482	sliceSize = ctx.moConfig.memSlice;	409	sliceSize = ctx.moConfig.memSlice;
483		410
484	- loopSrc[0].addr = src.addr;	411	+ loopSrc[0].addr_ = src.addr_;
485	loopSrc[0].token = src.token;	412	loopSrc[0].token = src.token;
486	- loopDst[0].addr = dst.addr;	413	+ loopDst[0].addr_ = dst.addr_;
487	loopDst[0].token = dst.token;	414	loopDst[0].token = dst.token;
488	loopLen[0] = sliceSize;	415	loopLen[0] = sliceSize;
489		416
490	loops.loopParam[0] = loopParam;	417	loops.loopParam[0] = loopParam;
491	- ccu::Variable paraCfg;	418	+ ccu::variable paraCfg;
492	paraCfg = GetParallelParam(ctx.moConfig.loopCount - 1, 0, 1);	419	paraCfg = GetParallelParam(ctx.moConfig.loopCount - 1, 0, 1);
493		420
494	- ccu::Variable offsetCfg;	421	+ ccu::variable offsetCfg;
495	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);	422	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);
496	- std::vector<ccu::Loop> grpLoops{*loops.loops[0]};	423	+ std::vector<ccu::loop> grpLoops{*loops.loops[0]};
497	- ccu::LoopGroup group(paraCfg, offsetCfg, ctx.moConfig.loopCount, grpLoops);	424	+ ccu::loop_group group(paraCfg, offsetCfg, ctx.moConfig.loopCount, grpLoops);
498	}	425	}
499		426
500	CCU_IF(goSize.parallelParam != 0)	427	CCU_IF(goSize.parallelParam != 0)
501	{	428	{
502	- src.addr += goSize.addrOffset;	429	+ src.addr_ += goSize.addrOffset;
503	- dst.addr += goSize.addrOffset;	430	+ dst.addr_ += goSize.addrOffset;
504		431
505	- loopSrc[0].addr = src.addr;	432	+ loopSrc[0].addr_ = src.addr_;
506	loopSrc[0].token = src.token;	433	loopSrc[0].token = src.token;
507	- loopDst[0].addr = dst.addr;	434	+ loopDst[0].addr_ = dst.addr_;
508	loopDst[0].token = dst.token;	435	loopDst[0].token = dst.token;
509	loopLen[0] = goSize.residual;	436	loopLen[0] = goSize.residual;
510		437
511	- src.addr += goSize.residual;	438	+ src.addr_ += goSize.residual;
512	- dst.addr += goSize.residual;	439	+ dst.addr_ += goSize.residual;
513		440
514	- ccu::Variable sliceSize;	441	+ ccu::variable sliceSize;
515	sliceSize = ctx.moConfig.memSlice;	442	sliceSize = ctx.moConfig.memSlice;
516		443
517	- loopSrc[1].addr = src.addr;	444	+ loopSrc[1].addr_ = src.addr_;
518	loopSrc[1].token = src.token;	445	loopSrc[1].token = src.token;
519	- loopDst[1].addr = dst.addr;	446	+ loopDst[1].addr_ = dst.addr_;
520	loopDst[1].token = dst.token;	447	loopDst[1].token = dst.token;
521	loopLen[1] = sliceSize;	448	loopLen[1] = sliceSize;
522		449
523	- ccu::Variable loopCfg0;	450	+ ccu::variable loopCfg0;
524	loopCfg0 = GetLoopParam(0, 0, 1);	451	loopCfg0 = GetLoopParam(0, 0, 1);
525	- ccu::Variable loopCfg1;	452	+ ccu::variable loopCfg1;
526	loopCfg1 = GetLoopParam(0, 0, 1);	453	loopCfg1 = GetLoopParam(0, 0, 1);
527	- ccu::Variable offsetCfg;	454	+ ccu::variable offsetCfg;
528	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);	455	offsetCfg = GetOffsetParam(ctx.moConfig.memSlice, ctx.moConfig.msInterleave, 1);
529		456
530	loops.loopParam[0] = loopCfg0;	457	loops.loopParam[0] = loopCfg0;
531	loops.loopParam[1] = loopCfg1;	458	loops.loopParam[1] = loopCfg1;
532	- std::vector<ccu::Loop> grpLoops{loops.loops[0], loops.loops[1]};	459	+ std::vector<ccu::loop> grpLoops{loops.loops[0], loops.loops[1]};
533	- ccu::LoopGroup group(goSize.parallelParam, offsetCfg, ctx.moConfig.loopCount, grpLoops);	460	+ ccu::loop_group group(goSize.parallelParam, offsetCfg, ctx.moConfig.loopCount, grpLoops);
534	}	461	}
535		462
536	return CCU_SUCCESS;	463	return CCU_SUCCESS;
	@@ -538,24 +465,24 @@ static CcuResult GroupCopy(
538		465
539	static CcuResult ExecuteAllGatherTransfer(AllGatherMesh1DMem2MemContext& ctx)	466	static CcuResult ExecuteAllGatherTransfer(AllGatherMesh1DMem2MemContext& ctx)
540	{	467	{
541	- ccu::LocalAddr src;	468	+ ccu::local_addr src;
542	- ccu::LocalAddr localDst;	469	+ ccu::local_addr localDst;
543	- std::vector<ccu::RemoteAddr> dst;	470	+ std::vector<ccu::remote_addr> dst;
544		471
545	dst.resize(ctx.arg->rankSize);	472	dst.resize(ctx.arg->rankSize);
546		473
547	- src.addr = ctx.input;	474	+ src.addr_ = ctx.input;
548	- src.addr += ctx.currentRankSliceInputOffset;	475	+ src.addr_ += ctx.currentRankSliceInputOffset;
549	src.token = ctx.token[ctx.arg->rankId];	476	src.token = ctx.token[ctx.arg->rankId];
550		477
551	for (uint32_t rankIdx = 0; rankIdx < ctx.arg->rankSize; rankIdx++) {	478	for (uint32_t rankIdx = 0; rankIdx < ctx.arg->rankSize; rankIdx++) {
552	if (rankIdx == ctx.arg->rankId) {	479	if (rankIdx == ctx.arg->rankId) {
553	- localDst.addr = ctx.output[rankIdx];	480	+ localDst.addr_ = ctx.output[rankIdx];
554	- localDst.addr += ctx.currentRankSliceOutputOffset;	481	+ localDst.addr_ += ctx.currentRankSliceOutputOffset;
555	localDst.token = ctx.token[rankIdx];	482	localDst.token = ctx.token[rankIdx];
556	} else {	483	} else {
557	- dst[rankIdx].addr = ctx.output[rankIdx];	484	+ dst[rankIdx].addr_ = ctx.output[rankIdx];
558	- dst[rankIdx].addr += ctx.currentRankSliceOutputOffset;	485	+ dst[rankIdx].addr_ += ctx.currentRankSliceOutputOffset;
559	dst[rankIdx].token = ctx.token[rankIdx];	486	dst[rankIdx].token = ctx.token[rankIdx];
560	}	487	}
561	}	488	}
	@@ -574,15 +501,15 @@ static CcuResult ExecuteAllGatherTransfer(AllGatherMesh1DMem2MemContext& ctx)
574	}	501	}
575		502
576	CCU_CHK_RET(GroupCopy(ctx, localDst, src, ctx.goSize));	503	CCU_CHK_RET(GroupCopy(ctx, localDst, src, ctx.goSize));
577	- CCU_CHK_RET(ccu::EventRecord(ctx.event, 1 << ctx.arg->rankId));	504	+ CCU_CHK_RET(ccu::event_record(ctx.event, 1 << ctx.arg->rankId));
578		505
579	const uint16_t totalMask = (1 << ctx.arg->rankSize) - 1;	506	const uint16_t totalMask = (1 << ctx.arg->rankSize) - 1;
580	- CCU_CHK_RET(ccu::EventWait(ctx.event, totalMask));	507	+ CCU_CHK_RET(ccu::event_wait(ctx.event, totalMask));
581		508
582	return CcuResult::CCU_SUCCESS;	509	return CcuResult::CCU_SUCCESS;
583	}	510	}
584		511
585	-__ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(CcuKernelArg arg)	512	+__ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(ccu_kernel_arg arg)
586	{	513	{
587	auto* kernelArg = static_cast<CcuKernelArgAllGatherMesh1DMem2Mem*>(arg);	514	auto* kernelArg = static_cast<CcuKernelArgAllGatherMesh1DMem2Mem*>(arg);
588		515
	@@ -599,16 +526,16 @@ __ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(CcuKernelArg arg)
599	uint32_t channelIdx = 0;	526	uint32_t channelIdx = 0;
600	for (uint64_t peerId = 0; peerId < ctx.arg->rankSize; peerId++) {	527	for (uint64_t peerId = 0; peerId < ctx.arg->rankSize; peerId++) {
601	if (peerId != ctx.arg->rankId) {	528	if (peerId != ctx.arg->rankId) {
602	- ctx.output[peerId] = ccu::GetResByChannel<ccu::Variable>(ctx.arg->channels[channelIdx], OUTPUT_XN_ID);	529	+ ctx.output[peerId] = ccu::GetResByChannel<ccu::variable>(ctx.arg->channels[channelIdx], OUTPUT_XN_ID);
603	- ctx.token[peerId] = ccu::GetResByChannel<ccu::Variable>(ctx.arg->channels[channelIdx], TOKEN_XN_ID);	530	+ ctx.token[peerId] = ccu::GetResByChannel<ccu::variable>(ctx.arg->channels[channelIdx], TOKEN_XN_ID);
604	channelIdx++;	531	channelIdx++;
605	}	532	}
606	}	533	}
607		534
608	- ccu::Array<ccu::Event> events(ctx.arg->eventHandle, ctx.arg->eventNum);	535	+ ccu::array<ccu::event> events(ctx.arg->eventHandle, ctx.arg->eventNum);
609	- ccu::EventWait(events[0]);	536	+ ccu::event_wait(events[0]);
610		537
611	- ccu::Array<ccu::Variable> vars(ctx.arg->varHandle, ctx.arg->varNum);	538	+ ccu::array<ccu::variable> vars(ctx.arg->varHandle, ctx.arg->varNum);
612	ctx.input = vars[0];	539	ctx.input = vars[0];
613	ctx.output[ctx.arg->rankId] = vars[1];	540	ctx.output[ctx.arg->rankId] = vars[1];
614	ctx.token[ctx.arg->rankId] = vars[2];	541	ctx.token[ctx.arg->rankId] = vars[2];
	@@ -621,26 +548,26 @@ __ccu_host__ void CcuAllGatherMesh1DMem2MemKernel(CcuKernelArg arg)
621	ctx.goSize.residual = vars[9];	548	ctx.goSize.residual = vars[9];
622		549
623	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	550	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
624	- CCU_CHK_VOID_RET(ccu::WriteVariableWithNotify(	551	+ CCU_CHK_VOID_RET(ccu::write_variable_with_notify(
625	ctx.arg->channels[i], ctx.output[ctx.arg->rankId], OUTPUT_XN_ID, CKE_IDX_0, 1 << OUTPUT_XN_ID));	552	ctx.arg->channels[i], ctx.output[ctx.arg->rankId], OUTPUT_XN_ID, CKE_IDX_0, 1 << OUTPUT_XN_ID));
626	- CCU_CHK_VOID_RET(ccu::WriteVariableWithNotify(	553	+ CCU_CHK_VOID_RET(ccu::write_variable_with_notify(
627	ctx.arg->channels[i], ctx.token[ctx.arg->rankId], TOKEN_XN_ID, CKE_IDX_0, 1 << TOKEN_XN_ID));	554	ctx.arg->channels[i], ctx.token[ctx.arg->rankId], TOKEN_XN_ID, CKE_IDX_0, 1 << TOKEN_XN_ID));
628	}	555	}
629		556
630	uint32_t allBit = (1 << OUTPUT_XN_ID) \| (1 << TOKEN_XN_ID);	557	uint32_t allBit = (1 << OUTPUT_XN_ID) \| (1 << TOKEN_XN_ID);
631	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	558	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
632	- CCU_CHK_VOID_RET(ccu::NotifyWait(ctx.arg->channels[i], CKE_IDX_0, allBit));	559	+ CCU_CHK_VOID_RET(ccu::notify_wait(ctx.arg->channels[i], CKE_IDX_0, allBit));
633	}	560	}
634		561
635	CCU_CHK_VOID_RET(ExecuteAllGatherTransfer(ctx));	562	CCU_CHK_VOID_RET(ExecuteAllGatherTransfer(ctx));
636		563
637	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	564	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
638	- CCU_CHK_VOID_RET(ccu::NotifyRecord(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));	565	+ CCU_CHK_VOID_RET(ccu::notify_record(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));
639	}	566	}
640	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {	567	for (uint32_t i = 0; i < ctx.arg->channelCount; i++) {
641	- CCU_CHK_VOID_RET(ccu::NotifyWait(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));	568	+ CCU_CHK_VOID_RET(ccu::notify_wait(ctx.arg->channels[i], CKE_IDX_0, 1 << POST_SYNC_ID));
642	}	569	}
643	- ccu::EventRecord(events[1]);	570	+ ccu::event_record(events[1]);
644	}	571	}
645		572
646	} // namespace ops_comm_ag	573	} // namespace ops_comm_ag
	@@ -756,27 +683,20 @@ HcclResult AllGatherAddCustom(
756	return HCCL_E_PARA;	683	return HCCL_E_PARA;
757	}	684	}
758	const uint8_t ccuDieId = static_cast<uint8_t>(dieId);	685	const uint8_t ccuDieId = static_cast<uint8_t>(dieId);
759	- const uint64_t dieMask = 1ULL << dieId;	686	+ CcuInsHandle insHandle{};
760	-	687	+ HCCLCHECK(asccomm_examples::QueryCcuInstance(comm, insHandle));
		tlq_2718 天前 QueryCcuInstance获取到的insHandle里为什么没有资源？
761	- CcuInsHandle insHandle{0};
762	- HCCLCHECK(CreateCcuInsWithEnabledDies(insHandle));
763	- HcclResult assignRet = HcclCommAssignCcuIns(comm, insHandle);
764	- if (assignRet != HCCL_SUCCESS) {
765	- (void)HcommCcuInsDestroy(insHandle);
766	- return assignRet;
767	- }
768		688
769	CcuMappedAddrs mappedAddrs{};	689	CcuMappedAddrs mappedAddrs{};
770	- CcuVariableHandle varHandle{0};	690	+ ccu_variable_handle varHandle{};
771	- HCCLCHECK(HcommCcuVariableAlloc(insHandle, ccuDieId, HCOMM_HOST_TASK_ARG_NUM, &varHandle));	691	+ HCCLCHECK(ConvertCcuToHccl(asccomm_ccu_variable_alloc(insHandle, ccuDieId, CCU_DIRECT_TASK_ARG_NUM, &varHandle)));
772	- for (uint32_t i = 0; i < HCOMM_HOST_TASK_ARG_NUM; i++) {	692	+ for (uint32_t i = 0; i < CCU_DIRECT_TASK_ARG_NUM; i++) {
773	- HCCLCHECK(HcommCcuVariableGetAddr(varHandle, i, &mappedAddrs.startVa[i]));	693	+ HCCLCHECK(ConvertCcuToHccl(asccomm_ccu_variable_get_addr(varHandle, i, &mappedAddrs.startVa[i])));
774	}	694	}
775		695
776	- CcuEventHandle eventHandle{0};	696	+ ccu_event_handle eventHandle{};
777	- HCCLCHECK(HcommCcuEventAlloc(insHandle, ccuDieId, CKE_NUM, &eventHandle));	697	+ HCCLCHECK(ConvertCcuToHccl(asccomm_ccu_event_alloc(insHandle, ccuDieId, CKE_NUM, &eventHandle)));
778	for (uint32_t i = 0; i < CKE_NUM; i++) {	698	for (uint32_t i = 0; i < CKE_NUM; i++) {
779	- HCCLCHECK(HcommCcuEventGetAddr(eventHandle, i, &mappedAddrs.eventVa[i]));	699	+ HCCLCHECK(ConvertCcuToHccl(asccomm_ccu_event_get_addr(eventHandle, i, &mappedAddrs.eventVa[i])));
780	}	700	}
781		701
782	auto kernelArg = std::make_shared<CcuKernelArgAllGatherMesh1DMem2Mem>();	702	auto kernelArg = std::make_shared<CcuKernelArgAllGatherMesh1DMem2Mem>();
	@@ -784,7 +704,7 @@ HcclResult AllGatherAddCustom(
784	kernelArg->rankId = rankId;	704	kernelArg->rankId = rankId;
785	kernelArg->channelCount = static_cast<uint32_t>(kernelChannels.size());	705	kernelArg->channelCount = static_cast<uint32_t>(kernelChannels.size());
786	kernelArg->varHandle = varHandle;	706	kernelArg->varHandle = varHandle;
787	- kernelArg->varNum = HCOMM_HOST_TASK_ARG_NUM;	707	+ kernelArg->varNum = CCU_DIRECT_TASK_ARG_NUM;
788	kernelArg->eventHandle = eventHandle;	708	kernelArg->eventHandle = eventHandle;
789	kernelArg->eventNum = CKE_NUM;	709	kernelArg->eventNum = CKE_NUM;
790	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	710	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -792,7 +712,7 @@ HcclResult AllGatherAddCustom(
792	}	712	}
793		713
794	uint64_t token = 0;	714	uint64_t token = 0;
795	- HCCLCHECK(HcommCcuGetMemToken(reinterpret_cast<uint64_t>(sendBuf), inputSize, &token));	715	+ HCCLCHECK(ConvertCcuToHccl(asccomm_ccu_get_mem_token(reinterpret_cast<uint64_t>(sendBuf), inputSize, &token)));
796		716
797	LoopGroupConfig config{};	717	LoopGroupConfig config{};
798	config.msInterleave = CCU_MS_INTERLEAVE;	718	config.msInterleave = CCU_MS_INTERLEAVE;
	@@ -800,7 +720,7 @@ HcclResult AllGatherAddCustom(
800	config.memSlice = CCU_MS_SIZE * CCU_LOCAL_COPY_MS_PER_LOOP;	720	config.memSlice = CCU_MS_SIZE * CCU_LOCAL_COPY_MS_PER_LOOP;
801	auto goSize = CalGoSize(inputSize, config);	721	auto goSize = CalGoSize(inputSize, config);
802		722
803	- uint64_t taskArgs[HCOMM_HOST_TASK_ARG_NUM] = {	723	+ uint64_t taskArgs[CCU_DIRECT_TASK_ARG_NUM] = {
804	reinterpret_cast<uint64_t>(sendBuf),	724	reinterpret_cast<uint64_t>(sendBuf),
805	reinterpret_cast<uint64_t>(recvBuf),	725	reinterpret_cast<uint64_t>(recvBuf),
806	token,	726	token,
	@@ -813,6 +733,7 @@ HcclResult AllGatherAddCustom(
813	goSize[3],	733	goSize[3],
814	};	734	};
815		735
		736	+ const uint64_t dieMask = 1ULL << dieId;
816	asccomm_ccu_schd schd = {1, 0, dieMask, 0};	737	asccomm_ccu_schd schd = {1, 0, dieMask, 0};
817	CcuAllGatherMesh1DMem2MemKernel<<<schd, insHandle, streamCcu>>>(kernelArg.get());	738	CcuAllGatherMesh1DMem2MemKernel<<<schd, insHandle, streamCcu>>>(kernelArg.get());
818		739

Mexamples/ccu/ccu_direct/03_add_allgather/CMakeLists.txt+5-5

	@@ -30,12 +30,10 @@ target_include_directories(demo PRIVATE
30	# CANN Toolkit头文件	30	# CANN Toolkit头文件
31	${ASCEND_CANN_PACKAGE_PATH}/include	31	${ASCEND_CANN_PACKAGE_PATH}/include
32	${ASCEND_CANN_PACKAGE_PATH}/include/hccl	32	${ASCEND_CANN_PACKAGE_PATH}/include/hccl
33	- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm	33	+ ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api
34	- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm/ccu
35	- ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api/ccu
36	${ASCEND_CANN_PACKAGE_PATH}/pkg_inc	34	${ASCEND_CANN_PACKAGE_PATH}/pkg_inc
37	- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/hcomm	35	+ # 样例公共辅助头（ccu_instance.h）
38	- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc/hcomm/ccu	36	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../common
39	)	37	)
40		38
41	target_compile_features(demo PRIVATE cxx_std_14)	39	target_compile_features(demo PRIVATE cxx_std_14)
	@@ -59,6 +57,7 @@ target_link_options(demo PRIVATE
59	-Wl,-z,relro	57	-Wl,-z,relro
60	-Wl,-z,now	58	-Wl,-z,now
61	-Wl,-z,noexecstack	59	-Wl,-z,noexecstack
		60	+ -Wl,-rpath,${ASCEND_CANN_PACKAGE_PATH}/lib64
62	)	61	)
63	target_link_directories(demo PRIVATE	62	target_link_directories(demo PRIVATE
64	${ASCEND_CANN_PACKAGE_PATH}/lib64	63	${ASCEND_CANN_PACKAGE_PATH}/lib64
	@@ -66,6 +65,7 @@ target_link_directories(demo PRIVATE
66	target_link_libraries(demo PRIVATE	65	target_link_libraries(demo PRIVATE
67	hcomm	66	hcomm
68	asccomm_ccu	67	asccomm_ccu
		68	+ asccomm_ccu_dataplane
69	c_sec	69	c_sec
70	ascendcl	70	ascendcl
71	acl_rt	71	acl_rt

Mexamples/ccu/ccu_direct/03_add_allgather/main.asc+105-187

	@@ -20,16 +20,13 @@
20	#include <hccl/hccl_comm.h>	20	#include <hccl/hccl_comm.h>
21	#include <hccl/hccl_rank_graph.h>	21	#include <hccl/hccl_rank_graph.h>
22	#include <hccl/hccl_res.h>	22	#include <hccl/hccl_res.h>
23	-#include <hccl/hccl_types.h>	23	+#include <hcomm/hcomm_primitives.h>
24	-#include <hccl/hcomm_primitives.h>	24	+#include <ccu/hcomm/ccu_api_types.h>
25	-#include <ccu/ccu_types.h>	25	+#include <ccu/hcomm/ccu_resource_api.h>
26	-#include <ccu/ccu_variable.hpp>	26	+#include <ccu/hcomm/ccu_primitives.hpp>
27	-#include <ccu/ccu_event.hpp>
28	-#include <ccu/ccu_primitives.hpp>
29		27
30	-#include "hccl_ccu_res.h"	28	+#include <ccu/ccu_host_launch.h>
31	-#include "ccu_host_launch.h"	29	+#include "../../../common/ccu_instance.h"
32	-#include "ccu_res.h"
33	#include "kernel_operator.h"	30	#include "kernel_operator.h"
34		31
35	#ifdef ACLCHECK	32	#ifdef ACLCHECK

Mexamples/ccu/ccu_reduce_scatter/op_host/alg_resource.cc+15-15

	@@ -22,33 +22,33 @@ namespace ops_hccl_rs {
22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;	22	constexpr uint32_t CHANNEL_NOTIFY_NUM = 3;
23		23
24	static HcclResult AllocThreadAndChannelResource(	24	static HcclResult AllocThreadAndChannelResource(
25	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,	25	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, std::vector<ChannelHandle>& kernelChannels,
26	uint32_t& kernelDie)	26	uint32_t& kernelDie)
27	{	27	{
28	ThreadHandle thread;	28	ThreadHandle thread;
29	constexpr uint32_t notifyNumOnMainThread = 0;	29	constexpr uint32_t notifyNumOnMainThread = 0;
30	RETURN_IF_HCCL_FAIL(	30	RETURN_IF_HCCL_FAIL(
31	- HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, param_.stream, notifyNumOnMainThread, &thread));	31	+ HcclThreadAcquireWithStream(comm, CommEngine::COMM_ENGINE_CCU, opParam.stream, notifyNumOnMainThread, &thread));
32	resCtxHost.threads.push_back(thread);	32	resCtxHost.threads.push_back(thread);
33		33
34	- if (param_.rankSize == 1) {	34	+ if (opParam.rankSize == 1) {
35	return HCCL_SUCCESS;	35	return HCCL_SUCCESS;
36	}	36	}
37		37
38	- uint32_t channel_num = param_.rankSize - 1;	38	+ uint32_t channel_num = opParam.rankSize - 1;
39	kernelChannels.resize(channel_num);	39	kernelChannels.resize(channel_num);
40		40
41	uint32_t channelIndex = 0;	41	uint32_t channelIndex = 0;
42	bool hasKernelDie = false;	42	bool hasKernelDie = false;
43	- for (uint32_t remoteRank = 0; remoteRank < param_.rankSize; remoteRank++) {	43	+ for (uint32_t remoteRank = 0; remoteRank < opParam.rankSize; remoteRank++) {
44	- if (remoteRank == param_.myRank) {	44	+ if (remoteRank == opParam.myRank) {
45	continue;	45	continue;
46	}	46	}
47		47
48	uint32_t netLayer = 0, listSize = 0;	48	uint32_t netLayer = 0, listSize = 0;
49	CommLink* linkList = nullptr;	49	CommLink* linkList = nullptr;
50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(	50	RETURN_IF_HCCL_FAIL(HcclRankGraphGetLinks(
51	- comm, netLayer, param_.myRank, remoteRank, &linkList,	51	+ comm, netLayer, opParam.myRank, remoteRank, &linkList,
52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息	52	&listSize)); // 获取 srcRank 和 dstRank 间的 link 信息
53		53
54	HcclChannelDesc desc;	54	HcclChannelDesc desc;
	@@ -69,7 +69,7 @@ static HcclResult AllocThreadAndChannelResource(
69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;	69	desc.remoteEndpoint.loc = link.dstEndpointDesc.loc;
70	uint32_t channelDie = 0;	70	uint32_t channelDie = 0;
71	RETURN_IF_HCCL_FAIL(	71	RETURN_IF_HCCL_FAIL(
72	- asccomm_examples::QueryEndpointDie(comm, param_.myRank, link.srcEndpointDesc, channelDie));	72	+ asccomm_examples::QueryEndpointDie(comm, opParam.myRank, link.srcEndpointDesc, channelDie));
73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));	73	RETURN_IF_HCCL_FAIL(asccomm_examples::MergeKernelDie(channelDie, hasKernelDie, kernelDie));
74	protocolExists = true;	74	protocolExists = true;
75	break;	75	break;
	@@ -88,7 +88,7 @@ static HcclResult AllocThreadAndChannelResource(
88	}	88	}
89		89
90	static HcclResult RegisterReduceScatterKernel(	90	static HcclResult RegisterReduceScatterKernel(
91	- HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,	91	+ HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost, const std::vector<ChannelHandle>& kernelChannels,
92	uint32_t kernelDie)	92	uint32_t kernelDie)
93	{	93	{
94	ccu_kernel_info kernelInfo;	94	ccu_kernel_info kernelInfo;
	@@ -100,8 +100,8 @@ static HcclResult RegisterReduceScatterKernel(
100	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuReduceScatterMesh1DMem2MemKernel);	100	kernelInfo.kernel_func = reinterpret_cast<void*>(CcuReduceScatterMesh1DMem2MemKernel);
101		101
102	auto kernel_arg = std::make_shared<CcuKernelArgReduceScatterMesh1DMem2Mem>();	102	auto kernel_arg = std::make_shared<CcuKernelArgReduceScatterMesh1DMem2Mem>();
103	- kernel_arg->rankSize = param_.rankSize;	103	+ kernel_arg->rankSize = opParam.rankSize;
104	- kernel_arg->rankId = param_.myRank;	104	+ kernel_arg->rankId = opParam.myRank;
105	kernelInfo.setKernelArg(kernel_arg);	105	kernelInfo.setKernelArg(kernel_arg);
106		106
107	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {	107	for (uint32_t i = 0; i < kernelChannels.size(); ++i) {
	@@ -140,17 +140,17 @@ static HcclResult RegisterReduceScatterKernel(
140	return HCCL_SUCCESS;	140	return HCCL_SUCCESS;
141	}	141	}
142		142
143	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost)	143	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost)
144	{	144	{
145	std::vector<ChannelHandle> kernelChannels;	145	std::vector<ChannelHandle> kernelChannels;
146	uint32_t kernelDie = 0;	146	uint32_t kernelDie = 0;
147	- RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, param_, resCtxHost, kernelChannels, kernelDie));	147	+ RETURN_IF_HCCL_FAIL(AllocThreadAndChannelResource(comm, opParam, resCtxHost, kernelChannels, kernelDie));
148		148
149	- if (param_.rankSize == 1) {	149	+ if (opParam.rankSize == 1) {
150	return HCCL_SUCCESS;	150	return HCCL_SUCCESS;
151	}	151	}
152		152
153	- RETURN_IF_HCCL_FAIL(RegisterReduceScatterKernel(comm, param_, resCtxHost, kernelChannels, kernelDie));	153	+ RETURN_IF_HCCL_FAIL(RegisterReduceScatterKernel(comm, opParam, resCtxHost, kernelChannels, kernelDie));
154	return HCCL_SUCCESS;	154	return HCCL_SUCCESS;
155	}	155	}
156		156

Mexamples/ccu/ccu_reduce_scatter/op_host/alg_resource.h+1-1

	@@ -38,7 +38,7 @@ struct OpParam {
38	uint32_t rankSize = 0;	38	uint32_t rankSize = 0;
39	};	39	};
40		40
41	-HcclResult AllocAlgResource(HcclComm comm, const OpParam& param_, AlgResourceCtx& resCtxHost);	41	+HcclResult AllocAlgResource(HcclComm comm, const OpParam& opParam, AlgResourceCtx& resCtxHost);
42		42
43	} // namespace ops_hccl_rs	43	} // namespace ops_hccl_rs
44		44

Mexamples/ccu/ccu_reduce_scatter/op_host/exec_op.cc+12-17

	@@ -12,7 +12,6 @@
12	#include <algorithm>	12	#include <algorithm>
13	#include <ccu/hcomm/ccu_api_types.h>	13	#include <ccu/hcomm/ccu_api_types.h>
14	#include <hccl/hcomm_primitives.h>	14	#include <hccl/hcomm_primitives.h>
15	-#include <hcomm/hcomm_ccu_launch_api.h>
16	#include "ccu/hcomm/ccu_launch.h"	15	#include "ccu/hcomm/ccu_launch.h"
17	#include "ccu/hcomm/ccu_resource_api.h"	16	#include "ccu/hcomm/ccu_resource_api.h"
18	#include "alg_resource.h"	17	#include "alg_resource.h"
	@@ -83,8 +82,8 @@ static std::vector<uint64_t> CalGoSize(uint64_t size, const loop_group_config& c
83	}	82	}
84		83
85	static HcclResult LaunchCcuKernelSlice(	84	static HcclResult LaunchCcuKernelSlice(
86	- const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t inputToken, uint64_t outputToken,	85	+ const OpParam& opParam, const AlgResourceCtx& resCtx, uint64_t inputAddr, uint64_t outputAddr, uint64_t inputToken,
87	- uint64_t rankSliceSize, uint64_t sliceCount, uint64_t dataTypeSize)	86	+ uint64_t outputToken, uint64_t rankSliceSize, uint64_t sliceCount, uint64_t dataTypeSize)
88	{	87	{
89	uint64_t sliceSize = sliceCount * dataTypeSize;	88	uint64_t sliceSize = sliceCount * dataTypeSize;
90		89
	@@ -99,12 +98,8 @@ static HcclResult LaunchCcuKernelSlice(
99	sliceSize, goSize[0], goSize[1], goSize[2], goSize[3],	98	sliceSize, goSize[0], goSize[1], goSize[2], goSize[3],
100	};	99	};
101		100
102	- HcommCcuLaunchContextPod launch_context{};
103	- if (HcommCcuGetLaunchContext(resCtx.threads[0], &launch_context) != HCCL_SUCCESS) {
104	- return HCCL_E_INTERNAL;
105	- }
106	CcuResult launchRet =	101	CcuResult launchRet =
107	- asccomm_ccu_kernel_launch(&launch_context, resCtx.ccuKernels[0], task_args.data(), task_args.size());	102	+ asccomm_ccu_kernel_launch(opParam.stream, resCtx.ccuKernels[0], task_args.data(), task_args.size());
108	if (launchRet != CCU_SUCCESS) {	103	if (launchRet != CCU_SUCCESS) {
109	return HCCL_E_INTERNAL;	104	return HCCL_E_INTERNAL;
110	}	105	}
	@@ -112,21 +107,21 @@ static HcclResult LaunchCcuKernelSlice(
112	return HCCL_SUCCESS;	107	return HCCL_SUCCESS;
113	}	108	}
114		109
115	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)	110	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx)
116	{	111	{
117	constexpr uint64_t dataTypeSize = sizeof(float);	112	constexpr uint64_t dataTypeSize = sizeof(float);
118	- const uint64_t rankSliceSize = param_.recvCount * dataTypeSize;	113	+ const uint64_t rankSliceSize = opParam.recvCount * dataTypeSize;
119	- const uint64_t inputSize = rankSliceSize * param_.rankSize;	114	+ const uint64_t inputSize = rankSliceSize * opParam.rankSize;
120	const uint64_t outputSize = rankSliceSize;	115	const uint64_t outputSize = rankSliceSize;
121	- const uint64_t count_ = param_.recvCount;	116	+ const uint64_t count_ = opParam.recvCount;
122		117
123	if (count_ == 0) { // 数据量为 0，直接返回	118	if (count_ == 0) { // 数据量为 0，直接返回
124	return HcclResult::HCCL_SUCCESS;	119	return HcclResult::HCCL_SUCCESS;
125	}	120	}
126		121
127	- if (param_.rankSize == 1) { // 单卡场景直接本地拷贝	122	+ if (opParam.rankSize == 1) { // 单卡场景直接本地拷贝
128	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(	123	RETURN_IF_HCCL_FAIL(static_cast<HcclResult>(
129	- HcommLocalCopyOnThread(resCtx.threads[0], param_.outputPtr, param_.inputPtr, outputSize)));	124	+ HcommLocalCopyOnThread(resCtx.threads[0], opParam.outputPtr, opParam.inputPtr, outputSize)));
130	return HCCL_SUCCESS;	125	return HCCL_SUCCESS;
131	}	126	}
132		127
	@@ -138,8 +133,8 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
138		133
139	uint64_t inputToken = 0;	134	uint64_t inputToken = 0;
140	uint64_t outputToken = 0;	135	uint64_t outputToken = 0;
141	- uint64_t baseInputAddr = reinterpret_cast<uint64_t>(param_.inputPtr);	136	+ uint64_t baseInputAddr = reinterpret_cast<uint64_t>(opParam.inputPtr);
142	- uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(param_.outputPtr);	137	+ uint64_t baseOutputAddr = reinterpret_cast<uint64_t>(opParam.outputPtr);
143	CcuResult tokenRet = asccomm_ccu_get_mem_token(baseInputAddr, inputSize, &inputToken);	138	CcuResult tokenRet = asccomm_ccu_get_mem_token(baseInputAddr, inputSize, &inputToken);
144	if (tokenRet != CCU_SUCCESS) {	139	if (tokenRet != CCU_SUCCESS) {
145	return HCCL_E_INTERNAL;	140	return HCCL_E_INTERNAL;
	@@ -155,7 +150,7 @@ HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx)
155	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;	150	uint64_t outputAddr = baseOutputAddr + processedDataCount * dataTypeSize;
156		151
157	RETURN_IF_HCCL_FAIL(LaunchCcuKernelSlice(	152	RETURN_IF_HCCL_FAIL(LaunchCcuKernelSlice(
158	- resCtx, inputAddr, outputAddr, inputToken, outputToken, rankSliceSize, sliceCount, dataTypeSize));	153	+ opParam, resCtx, inputAddr, outputAddr, inputToken, outputToken, rankSliceSize, sliceCount, dataTypeSize));
159		154
160	processedDataCount += sliceCount;	155	processedDataCount += sliceCount;
161	}	156	}

Mexamples/ccu/ccu_reduce_scatter/op_host/exec_op.h+1-1

	@@ -15,7 +15,7 @@
15		15
16	namespace ops_hccl_rs {	16	namespace ops_hccl_rs {
17		17
18	-HcclResult ExecOp(const OpParam& param_, const AlgResourceCtx& resCtx);	18	+HcclResult ExecOp(const OpParam& opParam, const AlgResourceCtx& resCtx);
19		19
20	} // namespace ops_hccl_rs	20	} // namespace ops_hccl_rs
21		21

Minclude/ccu/ccu_host_launch.h+3-31

 #include <stdint.h>
++#include "hcomm/hcomm_ccu_resource.h"
+-#if __has_include("hcomm/ccu/ccu_launch.h")
+-#include "hcomm/ccu/ccu_launch.h"
+-#define ASCCOMM_CCU_HOST_LAUNCH_HAS_HCOMM_TYPES 1
+-#elif __has_include("ccu/ccu_launch.h")
+-#include "ccu/ccu_launch.h"
+-#define ASCCOMM_CCU_HOST_LAUNCH_HAS_HCOMM_TYPES 1
+-#elif __has_include("ccu_launch.h")
+-#include "ccu_launch.h"
+-#define ASCCOMM_CCU_HOST_LAUNCH_HAS_HCOMM_TYPES 1
+-#endif
+-#endif
++#ifndef ASCCOMM_CCU_HOST_LAUNCH_HAS_ACLRT_STREAM
++typedef void* aclrtStream; // 无 acl 头环境的最小兜底，实际为 aclrtStream（void*）
+-    CCU_SUCCESS = 0,
+-    CCU_E_PARA = 1,
+-    CCU_E_PTR = 2,
+-    CCU_E_INTERNAL = 4,
+-    CCU_E_NOT_SUPPORT = 5,
+-    CCU_E_NOT_FOUND = 6,
+-    CCU_E_UNAVAIL = 7,
+-    CCU_E_RUNTIME = 15,
+-    CCU_E_DRV_START = 4096,
+-    CCU_E_DRV_INIT_FAILED = 4097,
+-    CCU_E_DRV_BUSY = 4098,
+-    CCU_E_DRV_END = 4224,
+-    CCU_E_RESERVED = 9216
+-} CcuResult;
+-
+-typedef uint64_t CcuInsHandle;
+-typedef void* aclrtStream;
 #endif
 #ifdef __cplusplus

Minclude/ccu/hcomm/ccu_launch.h+6-5

	@@ -12,16 +12,18 @@
12	#define CCU_LAUNCH_H	12	#define CCU_LAUNCH_H
13		13
14	// CCU kernel 注册与 launch 的 C ABI 头。	14	// CCU kernel 注册与 launch 的 C ABI 头。
15	-// 依赖的 Register Context 不透明句柄和 Launch Context POD 来自 hcomm 包内 header，	15	+// 依赖的 Register Context 不透明句柄来自 hcomm 包内 header，与 hccl 内部类型解耦，
16	-// 与 hccl 内部类型解耦，符合控制面/数据面 ABI 边界约束。	16	+// 符合控制面/数据面 ABI 边界约束。
17	#include "ccu/hcomm/ccu_api_types.h"	17	#include "ccu/hcomm/ccu_api_types.h"
18	-#include "hcomm/hcomm_ccu_launch.h"
19	#include "hcomm/hcomm_ccu_resource.h"	18	#include "hcomm/hcomm_ccu_resource.h"
		D D DragonBornHD8418 天前 [一般] 这个头现在写死要包含 hcomm/hcomm_ccu_resource.h，但这个文件只有 asc-comm 仓库自己留了一份：hcomm 安装包的 pkg_inc/hcomm/ 里没有它，asc-devkit 打包 asc-comm 头时又把 hcomm 目录整个排除了。结果就是用户拿 run 包编样例时，这里直接报"找不到头文件"。要么把它补进交付清单，要么恢复以前"找不到就用兜底定义"的写法，别让仓内的兜底头变成对外的硬依赖： `#if defined(__has_include) #if __has_include("hcomm/hcomm_ccu_resource.h") #include "hcomm/hcomm_ccu_resource.h" #define ASCCOMM_CCU_HAS_HCOMM_RESOURCE 1 #endif #endif #ifndef ASCCOMM_CCU_HAS_HCOMM_RESOURCE // <- 加这段兜底：没装 hcomm 头也能编 typedef uint64_t CcuInsHandle; // 定义与 hcomm_ccu_resource.h 保持一致 #endif`
20		19
21	#ifdef __cplusplus	20	#ifdef __cplusplus
22	extern "C" {	21	extern "C" {
23	#endif // __cplusplus	22	#endif // __cplusplus
24		23
		24	+// stream 直接以 void* 传递（调用方传 aclrtStream），避免公共头依赖 acl 头文件
		25	+typedef void* ccu_launch_stream;
		26	+
25	extern CcuResult asccomm_ccu_kernel_register_start(CcuInsHandle ins_handle);	27	extern CcuResult asccomm_ccu_kernel_register_start(CcuInsHandle ins_handle);
26		28
27	extern CcuResult asccomm_ccu_kernel_register(	29	extern CcuResult asccomm_ccu_kernel_register(
	@@ -33,8 +35,7 @@ extern CcuResult asccomm_ccu_kernel_register_end(CcuInsHandle ins_handle);
33	extern CcuResult asccomm_ccu_get_task_args_num(ccu_kernel_handle kernel_handle, uint32_t* task_args_num);	35	extern CcuResult asccomm_ccu_get_task_args_num(ccu_kernel_handle kernel_handle, uint32_t* task_args_num);
34		36
35	extern CcuResult asccomm_ccu_kernel_launch(	37	extern CcuResult asccomm_ccu_kernel_launch(
36	- const HcommCcuLaunchContextPod* launch_context, ccu_kernel_handle kernel_handle, const void* task_args,	38	+ ccu_launch_stream stream, ccu_kernel_handle kernel_handle, const void* task_args, uint32_t arg_num);
37	- uint32_t arg_num);
38		39
39	#ifdef __cplusplus	40	#ifdef __cplusplus
40	}	41	}

Minclude/hcomm/hcomm_types.h+5-0

	@@ -23,6 +23,9 @@
23	extern "C" {	23	extern "C" {
24	#endif	24	#endif
25		25
		26	+/* 已包含 hccl/hccl_types.h 时复用其稳定定义，避免与已安装 HCCL 包重定义枚举 */
		27	+#ifndef HCCL_TYPES_H_
		28	+
26	/** HCCL 接口返回码：0 成功，其余为错误码 */	29	/** HCCL 接口返回码：0 成功，其余为错误码 */
27	typedef enum {	30	typedef enum {
28	HCCL_SUCCESS = 0,	31	HCCL_SUCCESS = 0,
	@@ -88,6 +91,8 @@ typedef enum {
88	HCCL_DATA_TYPE_RESERVED = 255	91	HCCL_DATA_TYPE_RESERVED = 255
89	} HcclDataType;	92	} HcclDataType;
90		93
		94	+#endif // HCCL_TYPES_H_
		95	+
91	#ifdef __cplusplus	96	#ifdef __cplusplus
92	}	97	}
93	#endif	98	#endif

Msrc/ccu/CMakeLists.txt+7-5

	@@ -109,9 +109,8 @@ target_link_options(asccomm_ccu_dataplane PRIVATE
109	$<$<CONFIG:Release>:-Wl,--build-id=none>	109	$<$<CONFIG:Release>:-Wl,--build-id=none>
110	)	110	)
111		111
112	-# Keep the mainline host-launch wrapper as a separate compatibility library.	112	+# <<<>>> Host Launch 兼容层独立构建为 libasccomm_ccu.so，并通过数据面 C ABI
113	-# It uses the legacy HcommCcuKernel* registration ABI and must not be folded	113	+# 调用 libasccomm_ccu_dataplane.so 完成内核注册与下发。
114	-# into the data-plane SO, whose cross-SO interface is the POD-based API above.
115	add_library(asccomm_ccu_host_obj OBJECT	114	add_library(asccomm_ccu_host_obj OBJECT
116	${ASCCOMM_CCU_ROOT}/ccu_host_launch.cpp	115	${ASCCOMM_CCU_ROOT}/ccu_host_launch.cpp
117	)	116	)
	@@ -160,9 +159,12 @@ set_target_properties(asccomm_ccu PROPERTIES
160	OUTPUT_NAME asccomm_ccu	159	OUTPUT_NAME asccomm_ccu
161	LIBRARY_OUTPUT_DIRECTORY ${ASCCOMM_LIBRARY_OUTPUT_DIRECTORY}	160	LIBRARY_OUTPUT_DIRECTORY ${ASCCOMM_LIBRARY_OUTPUT_DIRECTORY}
162	)	161	)
163	-target_link_libraries(asccomm_ccu PRIVATE asccomm_ccu_headers)	162	+target_link_libraries(asccomm_ccu PRIVATE
164	-	163	+ asccomm_ccu_headers
		164	+ asccomm_ccu_dataplane
		165	+)
165	target_link_options(asccomm_ccu PRIVATE	166	target_link_options(asccomm_ccu PRIVATE
		167	+ -Wl,--no-undefined
166	-Wl,-z,relro -Wl,-z,now	168	-Wl,-z,relro -Wl,-z,now
167	$<$<CONFIG:Release>:-s>	169	$<$<CONFIG:Release>:-s>
168	$<$<CONFIG:Release>:-Wl,--build-id=none>	170	$<$<CONFIG:Release>:-Wl,--build-id=none>

Msrc/ccu/ccu_host_launch.cpp+37-140

	@@ -16,43 +16,18 @@
16	#include <mutex>	16	#include <mutex>
17	#include <unordered_map>	17	#include <unordered_map>
18		18
		19	+#include "ccu/hcomm/ccu_launch.h"
19	#include "xxhash_impl.inc"	20	#include "xxhash_impl.inc"
20		21
21	-#ifndef HCOMM_WEAK_SYMBOL
22	-#define HCOMM_WEAK_SYMBOL __attribute__((weak))
23	-#endif
24	-
25	-#ifndef ASCCOMM_CCU_HOST_LAUNCH_HAS_HCOMM_TYPES
26	-typedef uint64_t CcuKernelHandle;
27	-typedef void* CcuKernelArg;
28	-typedef uint64_t ThreadHandle;
29	-typedef int32_t CommEngine;
30	-typedef int32_t HcommResult;
31	-constexpr CommEngine COMM_ENGINE_CCU = static_cast<CommEngine>(0);
32	-#endif
33	-
34	-extern "C" {
35	-extern CcuResult HcommCcuKernelRegisterStart(CcuInsHandle insHandle) HCOMM_WEAK_SYMBOL;
36	-extern CcuResult HcommCcuKernelRegister(
37	- CcuInsHandle insHandle, uint32_t dieId, const char* kernelFuncName, const void* kernelFunc, const void** kernelArgs,
38	- uint32_t argNum, CcuKernelHandle* kernelHandle) HCOMM_WEAK_SYMBOL;
39	-extern CcuResult HcommCcuKernelRegisterEnd(CcuInsHandle insHandle) HCOMM_WEAK_SYMBOL;
40	-extern CcuResult HcommCcuGetTaskArgsNum(CcuKernelHandle kernelHandle, uint32_t* taskArgsNum) HCOMM_WEAK_SYMBOL;
41	-extern CcuResult HcommCcuKernelLaunch(
42	- ThreadHandle threadHandle, CcuKernelHandle kernelHandle, const void* taskArgs, uint32_t argNum) HCOMM_WEAK_SYMBOL;
43	-extern HcommResult HcommThreadAllocWithStream(
44	- CommEngine engine, aclrtStream stream, uint32_t notifyNum, ThreadHandle* thread);
45	-}
46	-
47	namespace {	22	namespace {
48	constexpr uint64_t CCU_DIE0_MASK = 0x01U;	23	constexpr uint64_t CCU_DIE0_MASK = 0x01U;
49	constexpr uint64_t CCU_DIE1_MASK = 0x02U;	24	constexpr uint64_t CCU_DIE1_MASK = 0x02U;
50	constexpr uint32_t CCU_DIE0_ID = 0U;	25	constexpr uint32_t CCU_DIE0_ID = 0U;
51	constexpr uint32_t CCU_DIE1_ID = 1U;	26	constexpr uint32_t CCU_DIE1_ID = 1U;
52	constexpr uint32_t CCU_SUPPORTED_NUM_BLOCKS = 1U;	27	constexpr uint32_t CCU_SUPPORTED_NUM_BLOCKS = 1U;
53	-constexpr uint32_t CCU_SQE_ARGS_LEN = 13U;
54	constexpr int32_t CCU_LOG_MODULE_ID = 5;	28	constexpr int32_t CCU_LOG_MODULE_ID = 5;
55		29
		30	+// Keep host launch behavior usable with test/minimal runtimes that omit the weak log backend.
56	#define ASCCOMM_CCU_LOG_ERROR(...) \	31	#define ASCCOMM_CCU_LOG_ERROR(...) \
57	do { \	32	do { \
58	if (DlogRecord != nullptr) { \	33	if (DlogRecord != nullptr) { \
	@@ -60,17 +35,6 @@ constexpr int32_t CCU_LOG_MODULE_ID = 5;
60	} \	35	} \
61	} while (0)	36	} while (0)
62		37
63	-bool IsCcuKernelLaunchApiAvailable()
64	-{
65	- auto registerStart = HcommCcuKernelRegisterStart;
66	- auto registerKernel = HcommCcuKernelRegister;
67	- auto registerEnd = HcommCcuKernelRegisterEnd;
68	- auto getTaskArgsNum = HcommCcuGetTaskArgsNum;
69	- auto kernelLaunch = HcommCcuKernelLaunch;
70	- return registerStart != nullptr && registerKernel != nullptr && registerEnd != nullptr &&
71	- getTaskArgsNum != nullptr && kernelLaunch != nullptr;
72	-}
73	-
74	bool IsSupportedSingleDieMask(uint64_t phyDieMask)	38	bool IsSupportedSingleDieMask(uint64_t phyDieMask)
75	{	39	{
76	return phyDieMask == CCU_DIE0_MASK \|\| phyDieMask == CCU_DIE1_MASK;	40	return phyDieMask == CCU_DIE0_MASK \|\| phyDieMask == CCU_DIE1_MASK;
	@@ -102,14 +66,17 @@ struct VoidKernelRegisterCtx {
102	void* packedArgs;	66	void* packedArgs;
103	};	67	};
104		68
105	-CcuResult VoidKernelTrampoline(CcuKernelArg arg);	69	+CcuResult VoidKernelTrampoline(ccu_kernel_arg arg);
106		70
107	CcuResult RegisterCcuKernel(	71	CcuResult RegisterCcuKernel(
108	const void* kernelFunc, const char* kernelName, const asccomm_launch_kernel_cfg* cfg, void* args,	72	const void* kernelFunc, const char* kernelName, const asccomm_launch_kernel_cfg* cfg, void* args,
109	- CcuKernelHandle& kernelHandle)	73	+ ccu_kernel_handle& kernelHandle)
110	{	74	{
111	- CcuResult ret = HcommCcuKernelRegisterStart(cfg->ccu_ins);	75	+ CcuResult ret = asccomm_ccu_kernel_register_start(cfg->ccu_ins);
112	if (ret != CCU_SUCCESS) {	76	if (ret != CCU_SUCCESS) {
		77	+ ASCCOMM_CCU_LOG_ERROR(
		78	+ "[%s] asccomm_ccu_kernel_register_start failed, ret[%d], ccuInsKey[0x%llx].", __func__, ret,
		79	+ static_cast<unsigned long long>(cfg->ccu_ins.ccuInsKey));
113	return ret;	80	return ret;
114	}	81	}
115		82
	@@ -123,22 +90,28 @@ CcuResult RegisterCcuKernel(
123		90
124	// kernelFunc do not support func that return void , only support return CcuResult	91	// kernelFunc do not support func that return void , only support return CcuResult
125	// kernelName非空时透传给注册侧(profiling名称等), 为空时注册侧使用默认名称	92	// kernelName非空时透传给注册侧(profiling名称等), 为空时注册侧使用默认名称
126	- ret = HcommCcuKernelRegister(	93	+ ret = asccomm_ccu_kernel_register(
127	cfg->ccu_ins, dieId, kernelName, reinterpret_cast<const void*>(VoidKernelTrampoline), kernelArgs, 1,	94	cfg->ccu_ins, dieId, kernelName, reinterpret_cast<const void*>(VoidKernelTrampoline), kernelArgs, 1,
128	&kernelHandle);	95	&kernelHandle);
129	if (ret != CCU_SUCCESS) {	96	if (ret != CCU_SUCCESS) {
130	- (void)HcommCcuKernelRegisterEnd(cfg->ccu_ins);	97	+ ASCCOMM_CCU_LOG_ERROR("[%s] asccomm_ccu_kernel_register failed, ret[%d], dieId[%u].", __func__, ret, dieId);
		98	+ (void)asccomm_ccu_kernel_register_end(cfg->ccu_ins);
131	return ret;	99	return ret;
132	}	100	}
133		101
134	- return HcommCcuKernelRegisterEnd(cfg->ccu_ins);	102	+ ret = asccomm_ccu_kernel_register_end(cfg->ccu_ins);
		103	+ if (ret != CCU_SUCCESS) {
		104	+ ASCCOMM_CCU_LOG_ERROR("[%s] asccomm_ccu_kernel_register_end failed, ret[%d].", __func__, ret);
		105	+ return ret;
		106	+ }
		107	+ return CCU_SUCCESS;
135	}	108	}
136		109
137	class KernelHandleCache {	110	class KernelHandleCache {
138	public:	111	public:
139	CcuResult Match(	112	CcuResult Match(
140	const void* kernelFunc, const char* kernelName, const asccomm_launch_kernel_cfg* cfg, void* args,	113	const void* kernelFunc, const char* kernelName, const asccomm_launch_kernel_cfg* cfg, void* args,
141	- CcuKernelHandle& kernelHandle, uint32_t& taskArgsNum)	114	+ ccu_kernel_handle& kernelHandle, uint32_t& taskArgsNum)
142	{	115	{
143	const uint64_t cacheTag = cfg->ccu_schd.binary_cache_tag;	116	const uint64_t cacheTag = cfg->ccu_schd.binary_cache_tag;
144	if (cacheTag == 0U) {	117	if (cacheTag == 0U) {
	@@ -149,7 +122,7 @@ public:
149	auto iter = cache_.find(cacheTag);	122	auto iter = cache_.find(cacheTag);
150	if (iter != cache_.end()) {	123	if (iter != cache_.end()) {
151	kernelHandle = iter->second;	124	kernelHandle = iter->second;
152	- return HcommCcuGetTaskArgsNum(kernelHandle, &taskArgsNum);	125	+ return asccomm_ccu_get_task_args_num(kernelHandle, &taskArgsNum);
153	}	126	}
154		127
155	CcuResult ret = RegisterAndGetTaskArgsNum(kernelFunc, kernelName, cfg, args, kernelHandle, taskArgsNum);	128	CcuResult ret = RegisterAndGetTaskArgsNum(kernelFunc, kernelName, cfg, args, kernelHandle, taskArgsNum);
	@@ -163,21 +136,24 @@ public:
163	private:	136	private:
164	CcuResult RegisterAndGetTaskArgsNum(	137	CcuResult RegisterAndGetTaskArgsNum(
165	const void* kernelFunc, const char* kernelName, const asccomm_launch_kernel_cfg* cfg, void* args,	138	const void* kernelFunc, const char* kernelName, const asccomm_launch_kernel_cfg* cfg, void* args,
166	- CcuKernelHandle& kernelHandle, uint32_t& taskArgsNum)	139	+ ccu_kernel_handle& kernelHandle, uint32_t& taskArgsNum)
167	{	140	{
168	CcuResult ret = RegisterCcuKernel(kernelFunc, kernelName, cfg, args, kernelHandle);	141	CcuResult ret = RegisterCcuKernel(kernelFunc, kernelName, cfg, args, kernelHandle);
169	if (ret != CCU_SUCCESS) {	142	if (ret != CCU_SUCCESS) {
170	return ret;	143	return ret;
171	}	144	}
172	- ret = HcommCcuGetTaskArgsNum(kernelHandle, &taskArgsNum);	145	+ ret = asccomm_ccu_get_task_args_num(kernelHandle, &taskArgsNum);
173	if (ret != CCU_SUCCESS) {	146	if (ret != CCU_SUCCESS) {
		147	+ ASCCOMM_CCU_LOG_ERROR(
		148	+ "[%s] asccomm_ccu_get_task_args_num failed, ret[%d], kernelHandle[0x%llx].", __func__, ret,
		149	+ static_cast<unsigned long long>(kernelHandle));
174	return ret;	150	return ret;
175	}	151	}
176	return CCU_SUCCESS;	152	return CCU_SUCCESS;
177	}	153	}
178		154
179	std::mutex mutex_;	155	std::mutex mutex_;
180	- std::unordered_map<uint64_t, CcuKernelHandle> cache_;	156	+ std::unordered_map<uint64_t, ccu_kernel_handle> cache_;
181	};	157	};
182		158
183	KernelHandleCache& GetKernelHandleCache()	159	KernelHandleCache& GetKernelHandleCache()
	@@ -186,7 +162,7 @@ KernelHandleCache& GetKernelHandleCache()
186	return cache;	162	return cache;
187	}	163	}
188		164
189	-CcuResult VoidKernelTrampoline(CcuKernelArg arg)	165	+CcuResult VoidKernelTrampoline(ccu_kernel_arg arg)
190	{	166	{
191	auto* ctx = static_cast<VoidKernelRegisterCtx*>(arg);	167	auto* ctx = static_cast<VoidKernelRegisterCtx*>(arg);
192	if (ctx == nullptr \|\| ctx->kernelFunc == nullptr \|\| ctx->packedArgs == nullptr) {	168	if (ctx == nullptr \|\| ctx->kernelFunc == nullptr \|\| ctx->packedArgs == nullptr) {
	@@ -198,91 +174,6 @@ CcuResult VoidKernelTrampoline(CcuKernelArg arg)
198	return CCU_SUCCESS;	174	return CCU_SUCCESS;
199	}	175	}
200		176
201	-CcuResult ConvertHcommResult(HcommResult ret)
202	-{
203	- switch (ret) {
204	- case 0:
205	- return CCU_SUCCESS;
206	- case 1:
207	- return CCU_E_PARA;
208	- case 2:
209	- return CCU_E_PTR;
210	- case 3:
211	- return CCU_E_INTERNAL;
212	- case 4:
213	- return CCU_E_INTERNAL;
214	- case 5:
215	- return CCU_E_NOT_SUPPORT;
216	- case 6:
217	- return CCU_E_NOT_FOUND;
218	- case 7:
219	- return CCU_E_UNAVAIL;
220	- default:
221	- return CCU_E_RUNTIME;
222	- }
223	-}
224	-
225	-class ThreadHandleCache {
226	-public:
227	- CcuResult GetOrCreate(aclrtStream stream, ThreadHandle& thread)
228	- {
229	- std::lock_guard<std::mutex> lock(mutex_);
230	- auto iter = cache_.find(stream);
231	- if (iter != cache_.end()) {
232	- thread = iter->second;
233	- return CCU_SUCCESS;
234	- }
235	-
236	- ThreadHandle newThread = 0;
237	- const HcommResult allocRet = HcommThreadAllocWithStream(COMM_ENGINE_CCU, stream, 0, &newThread);
238	- if (allocRet != 0) {
239	- ASCCOMM_CCU_LOG_ERROR(
240	- "[%s] HcommThreadAllocWithStream failed, ret[%d], stream[%p].", __func__, allocRet, stream);
241	- return ConvertHcommResult(allocRet);
242	- }
243	- if (newThread == 0) {
244	- ASCCOMM_CCU_LOG_ERROR("[%s] HcommThreadAllocWithStream returned an empty handle.", __func__);
245	- return CCU_E_INTERNAL;
246	- }
247	-
248	- cache_.emplace(stream, newThread);
249	- thread = newThread;
250	- return CCU_SUCCESS;
251	- }
252	-
253	-private:
254	- std::mutex mutex_;
255	- std::unordered_map<aclrtStream, ThreadHandle> cache_;
256	-};
257	-
258	-ThreadHandleCache& GetThreadHandleCache()
259	-{
260	- static ThreadHandleCache cache;
261	- return cache;
262	-}
263	-
264	-CcuResult LaunchCcuKernelWithStream(
265	- CcuKernelHandle kernelHandle, const void* taskArgs, uint32_t taskArgsNum, aclrtStream stream)
266	-{
267	- if (stream == nullptr) {
268	- ASCCOMM_CCU_LOG_ERROR("[%s] stream is nullptr.", __func__);
269	- return CCU_E_PTR;
270	- }
271	- if (taskArgsNum > CCU_SQE_ARGS_LEN) {
272	- ASCCOMM_CCU_LOG_ERROR(
273	- "[%s] task argument count[%u] exceeds maximum[%u].", __func__, taskArgsNum, CCU_SQE_ARGS_LEN);
274	- return CCU_E_NOT_SUPPORT;
275	- }
276	-
277	- ThreadHandle threadHandle = 0;
278	- CcuResult ret = GetThreadHandleCache().GetOrCreate(stream, threadHandle);
279	- if (ret != CCU_SUCCESS) {
280	- return ret;
281	- }
282	-
283	- return HcommCcuKernelLaunch(threadHandle, kernelHandle, taskArgs, taskArgsNum);
284	-}
285	-
286	} // namespace	177	} // namespace
287		178
288	extern "C" uint64_t asccomm_ccu_get_launch_hash_tag(const char* tag)	179	extern "C" uint64_t asccomm_ccu_get_launch_hash_tag(const char* tag)
	@@ -297,24 +188,30 @@ extern "C" ccu_result asccomm_ccu_host_kernel_launch(
297	const void* kernel_func, const char* kernel_name, const asccomm_launch_kernel_cfg* cfg, void* args)	188	const void* kernel_func, const char* kernel_name, const asccomm_launch_kernel_cfg* cfg, void* args)
298	{	189	{
299	if (kernel_func == nullptr \|\| cfg == nullptr \|\| args == nullptr) {	190	if (kernel_func == nullptr \|\| cfg == nullptr \|\| args == nullptr) {
		191	+ ASCCOMM_CCU_LOG_ERROR(
		192	+ "[%s] invalid input, kernelFunc[%p], cfg[%p], args[%p].", __func__, kernel_func, cfg, args);
300	return CCU_E_PTR;	193	return CCU_E_PTR;
301	}	194	}
302	CcuResult ret = ValidateLaunchCfg(cfg);	195	CcuResult ret = ValidateLaunchCfg(cfg);
303	if (ret != CCU_SUCCESS) {	196	if (ret != CCU_SUCCESS) {
304	return ret;	197	return ret;
305	}	198	}
306	- if (!IsCcuKernelLaunchApiAvailable()) {
307	- return CCU_E_NOT_SUPPORT;
308	- }
309		199
310	- CcuKernelHandle kernelHandle = 0;	200	+ ccu_kernel_handle kernelHandle = 0;
311	uint32_t taskArgsNum = 0;	201	uint32_t taskArgsNum = 0;
312	ret = GetKernelHandleCache().Match(kernel_func, kernel_name, cfg, args, kernelHandle, taskArgsNum);	202	ret = GetKernelHandleCache().Match(kernel_func, kernel_name, cfg, args, kernelHandle, taskArgsNum);
313	if (ret != CCU_SUCCESS) {	203	if (ret != CCU_SUCCESS) {
314	return ret;	204	return ret;
315	}	205	}
316		206
317	- return LaunchCcuKernelWithStream(kernelHandle, args, taskArgsNum, cfg->stream);	207	+ ret = asccomm_ccu_kernel_launch(cfg->stream, kernelHandle, args, taskArgsNum);
		208	+ if (ret != CCU_SUCCESS) {
		209	+ ASCCOMM_CCU_LOG_ERROR(
		210	+ "[%s] asccomm_ccu_kernel_launch failed, ret[%d], kernelHandle[0x%llx], stream[%p].", __func__, ret,
		211	+ static_cast<unsigned long long>(kernelHandle), cfg->stream);
		212	+ return ret;
		213	+ }
		214	+ return CCU_SUCCESS;
318	}	215	}
319		216
320	extern "C" CcuResult HcommCcuHostKernelLaunch(	217	extern "C" CcuResult HcommCcuHostKernelLaunch(

Msrc/ccu/hcomm/primitives/ccu_launch.cc+39-192

	@@ -10,9 +10,7 @@
10		10
11	#include "ccu/hcomm/ccu_launch.h"	11	#include "ccu/hcomm/ccu_launch.h"
12	#include "ccu/hcomm/ccu_resource_api.h"	12	#include "ccu/hcomm/ccu_resource_api.h"
13	-#include "hcomm/hcomm_ccu_launch.h"
14		13
15	-#include <chrono>
16	#include <vector>	14	#include <vector>
17		15
18	// 同步 launch 需要 acl 运行时接口（rt 相关 API），与 runtime 头配合使用	16	// 同步 launch 需要 acl 运行时接口（rt 相关 API），与 runtime 头配合使用
	@@ -27,9 +25,10 @@
27	#include "hcomm/resource/kernel/ccu_kernel_registry_mgr.h"	25	#include "hcomm/resource/kernel/ccu_kernel_registry_mgr.h"
28	#include "hcomm/resource/common/ccu_var_event_res_mgr.h"	26	#include "hcomm/resource/common/ccu_var_event_res_mgr.h"
29		27
30	-#include "hcomm/resource/microcode/ccu_assist_v1.h"	28	+// 在不进入内核注册流程的情况下准备单实例资源快照。
31	-	29	+// 变量/事件分配接口可能先于 <<<>>> 调用，因此需要执行此操作。
32	-CcuResult asccomm_ccu_kernel_register_start(CcuInsHandle ins_handle)	30	+static CcuResult ccu_prepare_instance(
		31	+ CcuInsHandle ins_handle, asc::ccu_kernel_registry*& ccu_ins, int32_t& dev_logic_id)
33	{	32	{
34	if (ins_handle.ccuInsKey == 0 \|\| ins_handle.ccuInsPtr == nullptr) {	33	if (ins_handle.ccuInsKey == 0 \|\| ins_handle.ccuInsPtr == nullptr) {
35	HCCL_ERROR(	34	HCCL_ERROR(
	@@ -37,31 +36,36 @@ CcuResult asccomm_ccu_kernel_register_start(CcuInsHandle ins_handle)
37	static_cast<unsigned long long>(ins_handle.ccuInsKey), static_cast<const void*>(ins_handle.ccuInsPtr));	36	static_cast<unsigned long long>(ins_handle.ccuInsKey), static_cast<const void*>(ins_handle.ccuInsPtr));
38	return CcuResult::CCU_E_PARA;	37	return CcuResult::CCU_E_PARA;
39	}	38	}
40	- const auto* instance = ins_handle.ccuInsPtr;	39	+
41	- // Instance 不再携带设备号：以当前线程设备定位 per-device registry/kernel mgr	40	+ dev_logic_id = asc::get_current_ccu_device_logic_id();
42	- const int32_t dev_logic_id = asc::get_current_ccu_device_logic_id();
43	if (dev_logic_id < 0) {	41	if (dev_logic_id < 0) {
44	HCCL_ERROR("[%s] failed, current thread has no device set.", __func__);	42	HCCL_ERROR("[%s] failed, current thread has no device set.", __func__);
45	return CcuResult::CCU_E_UNAVAIL;	43	return CcuResult::CCU_E_UNAVAIL;
46	}	44	}
47	- asc::ccu_kernel_registry* ccu_ins = nullptr;	45	+
48	CCU_CHK_RET(	46	CCU_CHK_RET(
49	asc::ccu_kernel_registry_mgr::get_instance(dev_logic_id).get_or_create(dev_logic_id, ins_handle, ccu_ins));	47	asc::ccu_kernel_registry_mgr::get_instance(dev_logic_id).get_or_create(dev_logic_id, ins_handle, ccu_ins));
50	CCU_CHK_PTR_NULL(ccu_ins);	48	CCU_CHK_PTR_NULL(ccu_ins);
		49	+ return ccu_ins->load_register_context(*ins_handle.ccuInsPtr);
		50	+}
		51	+
		52	+CcuResult asccomm_ccu_kernel_register_start(CcuInsHandle ins_handle)
		53	+{
		54	+ // Instance 不再携带设备号：以当前线程设备定位 per-device registry/kernel mgr
		55	+ int32_t dev_logic_id = -1;
		56	+ asc::ccu_kernel_registry* ccu_ins = nullptr;
		57	+ CCU_CHK_RET(ccu_prepare_instance(ins_handle, ccu_ins, dev_logic_id));
51	CCU_CHK_RET(ccu_ins->begin_register());	58	CCU_CHK_RET(ccu_ins->begin_register());
52		59
53	- CcuResult ret = ccu_ins->load_register_context(*instance);	60	+ auto* res_snapshot = ccu_ins->get_res_snapshot();
54	- if (ret == CcuResult::CCU_SUCCESS) {	61	+ CCU_CHK_PTR_NULL(res_snapshot);
55	- auto* res_snapshot = ccu_ins->get_res_snapshot();	62	+ CcuResult ret = asc::ccu_kernel_mgr::get_instance(dev_logic_id).configure(*res_snapshot);
56	- CCU_CHK_PTR_NULL(res_snapshot);
57	- ret = asc::ccu_kernel_mgr::get_instance(dev_logic_id).configure(*res_snapshot);
58	- }
59	if (ret != CcuResult::CCU_SUCCESS) {	63	if (ret != CcuResult::CCU_SUCCESS) {
60	(void)ccu_ins->end_register();	64	(void)ccu_ins->end_register();
61	HCCL_ERROR("[%s] failed to load register context, ret[%d].", __func__, ret);	65	HCCL_ERROR("[%s] failed to load register context, ret[%d].", __func__, ret);
62	return ret;	66	return ret;
63	}	67	}
64	- asc::set_current_ccu_control_ops(instance->ascCustom);	68	+ asc::set_current_ccu_control_ops(ins_handle.ccuInsPtr->ascCustom);
65	return CcuResult::CCU_SUCCESS;	69	return CcuResult::CCU_SUCCESS;
66	}	70	}
67		71
	@@ -84,7 +88,6 @@ CcuResult asccomm_ccu_kernel_register(
84	const void** kernel_args, uint32_t arg_num, ccu_kernel_handle* kernel_handle)	88	const void** kernel_args, uint32_t arg_num, ccu_kernel_handle* kernel_handle)
85	{	89	{
86	HCCL_RUN_INFO("Entry-%s", __func__);	90	HCCL_RUN_INFO("Entry-%s", __func__);
87	- const auto start_time = std::chrono::steady_clock::now();
88		91
89	CCU_CHK_PTR_NULL(kernel_func);	92	CCU_CHK_PTR_NULL(kernel_func);
90	CCU_CHK_PTR_NULL(kernel_handle);	93	CCU_CHK_PTR_NULL(kernel_handle);
	@@ -118,9 +121,6 @@ CcuResult asccomm_ccu_kernel_register(
118	}	121	}
119		122
120	*kernel_handle = new_handle;	123	*kernel_handle = new_handle;
121	- const auto duration =
122	- std::chrono::duration_cast<std::chrono::microseconds>(std::chrono::steady_clock::now() - start_time);
123	- HCCL_INFO("[%s] success, take time [%lld]us.", __func__, duration.count());
124	return CcuResult::CCU_SUCCESS;	124	return CcuResult::CCU_SUCCESS;
125	}	125	}
126		126
	@@ -158,116 +158,10 @@ CcuResult asccomm_ccu_get_task_args_num(ccu_kernel_handle kernel_handle, uint32_
158	return CcuResult::CCU_SUCCESS;	158	return CcuResult::CCU_SUCCESS;
159	}	159	}
160		160
161	-static CcuResult construct_ccu_detail_info(	161	+// notifywait 默认等待时长，27*68=1836，沿用 hcomm NOTIFY_DEFAULT_WAIT_TIME 的 notifywait 默认等待时长
162	- const std::vector<asc::ccu_profiling_info>& all_ccu_profiling_info, bool is_save_profiling_data,	162	+constexpr uint32_t CCU_NOTIFY_DEFAULT_WAIT_TIME = 27U * 68U;
163	- std::vector<HcommCcuProfileDetailPod>& converted)
164	-{
165	- if (all_ccu_profiling_info.empty() \|\| !is_save_profiling_data) {
166	- return CcuResult::CCU_SUCCESS;
167	- }
168		163
169	- converted.resize(all_ccu_profiling_info.size());	164	+static CcuResult launch_ccu_tasks(const asc::ccu_task_param& param, const aclrtStream stream)
170	- // profiling 明细跨动态库时只传递 POD，hcomm 在 callback 前再还原为原有 C++ 类型。
171	- for (uint32_t idx = 0; idx < all_ccu_profiling_info.size(); ++idx) {
172	- const auto& src = all_ccu_profiling_info[idx];
173	- auto& dst = converted[idx];
174	- if (src.name.size() >= HCOMM_CCU_PROFILE_NAME_CAPACITY) {
175	- HCCL_ERROR(
176	- "[%s] profiling name length[%zu] exceeds maximum[%u].", __func__, src.name.size(),
177	- HCOMM_CCU_PROFILE_NAME_CAPACITY - 1U);
178	- return CcuResult::CCU_E_PARA;
179	- }
180	- dst.header = {HCOMM_CCU_LAUNCH_ABI_VERSION, HCOMM_CCU_PROFILE_DETAIL_MAGIC_WORD, sizeof(dst), 0};
181	- dst.nameLength = static_cast<uint32_t>(src.name.size());
182	- (void)memcpy_s(dst.name, sizeof(dst.name), src.name.data(), src.name.size());
183	- dst.profilingType = src.type;
184	- dst.dieId = src.die_id;
185	- dst.missionId = src.mission_id;
186	- dst.instructionId = src.instr_id;
187	- dst.reduceOpType = src.reduce_op_type;
188	- dst.inputDataType = src.input_data_type;
189	- dst.outputDataType = src.output_data_type;
190	- dst.dataSize = src.data_size;
191	- dst.ckeId = src.cke_id;
192	- dst.mask = src.mask;
193	- (void)memcpy_s(dst.channelId, sizeof(dst.channelId), src.channel_id, sizeof(src.channel_id));
194	- (void)memcpy_s(dst.remoteRankId, sizeof(dst.remoteRankId), src.remote_rank_id, sizeof(src.remote_rank_id));
195	- (void)memcpy_s(dst.channelHandle, sizeof(dst.channelHandle), src.channel_handle, sizeof(src.channel_handle));
196	- }
197	- return CcuResult::CCU_SUCCESS;
198	-}
199	-
200	-static HcommCcuTaskProfilePod construct_ccu_task_profile(
201	- const asc::ccu_task_param& ccu_param, const ccu_kernel_handle kernel_handle, uint64_t begin_time, uint64_t end_time,
202	- bool is_master)
203	-{
204	- // POD 头部固定携带版本、类型和大小，供 hcomm 在解释后续 profiling 字段前校验 ABI。
205	- HcommCcuTaskProfilePod task{};
206	- task.header = {HCOMM_CCU_LAUNCH_ABI_VERSION, HCOMM_CCU_TASK_PROFILE_MAGIC_WORD, sizeof(task), 0};
207	- task.beginCycle = begin_time;
208	- task.endCycle = end_time;
209	- task.dieId = ccu_param.die_id;
210	- task.missionId = ccu_param.mission_id;
211	- task.instructionId = ccu_param.inst_start_id;
212	- task.isMaster = is_master ? 1U : 0U;
213	- task.kernelHandle = kernel_handle;
214	- task.diagnose = asc::asccomm_ccu_diagnose;
215	- task.reserved = 0;
216	- return task;
217	-}
218	-
219	-static void log_ccu_task_info(
220	- const std::vector<asc::ccu_task_param>& ccu_params, const ccu_kernel_handle kernel_handle,
221	- uint32_t exec_time_out_sec)
222	-{
223	- if (CheckLogLevel(asccomm_ccu_log_module_id, DLOG_INFO) != 1) {
224	- return;
225	- }
226	- for (uint32_t idx = 0; idx < ccu_params.size(); idx++) {
227	- const auto& param = ccu_params[idx];
228	- HCCL_INFO(
229	- "[%s] start ccu task, dieId[%u], missionId[%u], execMissionId[%u], instStartId[%u], instCnt[%u], "
230	- "argSize[%u], timeout[%u]s, executeId[0x%llx], ccuKernelHandle[0x%llx]",
231	- __func__, param.die_id, param.mission_id, param.mission_id, param.inst_start_id, param.inst_cnt,
232	- param.arg_size, exec_time_out_sec, kernel_handle, kernel_handle);
233	- }
234	-}
235	-
236	-static void construct_profiling_info_log(const std::vector<asc::ccu_profiling_info>& all_ccu_profiling_info)
237	-{
238	- if (CheckLogLevel(asccomm_ccu_log_module_id, DLOG_INFO) != 1) {
239	- return;
240	- }
241	- for (const asc::ccu_profiling_info& prof_info : all_ccu_profiling_info) {
242	- for (int idx = 0; idx < asc::ccu_max_channel_num; idx++) {
243	- if (prof_info.channel_id[idx] == asc::invalid_value_channelid) {
244	- break;
245	- }
246	- HCCL_INFO(
247	- "[%s]idx[%d]: channelId[%u], channelHandle[0x%llx]", __func__, idx, prof_info.channel_id[idx],
248	- prof_info.channel_handle[idx]);
249	- }
250	- }
251	-}
252	-
253	-static CcuResult construct_profiling_info(
254	- asc::ccu_kernel* kernel, const uint64_t* task_args, uint32_t arg_num,
255	- std::vector<asc::ccu_profiling_info>& all_ccu_profiling_info, bool is_save_profiling_data)
256	-{
257	- if (!is_save_profiling_data) {
258	- return CcuResult::CCU_SUCCESS;
259	- }
260	-
261	- CCU_CHK_RET(kernel->get_ccu_profiling_info(task_args, arg_num, all_ccu_profiling_info));
262	- if (all_ccu_profiling_info.empty()) {
263	- return CcuResult::CCU_SUCCESS;
264	- }
265	- construct_profiling_info_log(all_ccu_profiling_info);
266	- return CcuResult::CCU_SUCCESS;
267	-}
268	-
269	-static CcuResult launch_ccu_tasks(
270	- const asc::ccu_task_param& param, const aclrtStream stream, uint32_t exec_time_out_sec)
271	{	165	{
272	rtCcuTaskInfo_t task_info{};	166	rtCcuTaskInfo_t task_info{};
273	task_info.dieId = param.die_id;	167	task_info.dieId = param.die_id;
	@@ -276,7 +170,7 @@ static CcuResult launch_ccu_tasks(
276	task_info.instCnt = param.inst_cnt;	170	task_info.instCnt = param.inst_cnt;
277	task_info.key = param.key;	171	task_info.key = param.key;
278	task_info.argSize = param.arg_size;	172	task_info.argSize = param.arg_size;
279	- task_info.timeout = exec_time_out_sec;	173	+ task_info.timeout = CCU_NOTIFY_DEFAULT_WAIT_TIME;
		D D DragonBornHD8418 天前 [一般] 超时时间从"可以配"变成"写死"了：原来这个值由 hcomm 侧环境变量（EnvConfig 的 GetExecTimeOut）控制，现场有靠它调大超时的长任务，这次会悄悄失效、任务可能被提前判超时。至少把这个魔数的来历写明白，要保留可配能力就给 launch 留个可选参数： `// 2768=1836，沿用 hcomm NOTIFY_DEFAULT_WAIT_TIME 的 notifywait 默认等待时长 constexpr uint32_t CCU_NOTIFY_DEFAULT_WAIT_TIME = 27U 68U; // <- 补上出处注释`
280	std::copy(std::begin(param.args), std::end(param.args), std::begin(task_info.args));	174	std::copy(std::begin(param.args), std::end(param.args), std::begin(task_info.args));
281		175
282	auto ret = rtCCULaunch(&task_info, stream);	176	auto ret = rtCCULaunch(&task_info, stream);
	@@ -288,89 +182,42 @@ static CcuResult launch_ccu_tasks(
288	return CcuResult::CCU_SUCCESS;	182	return CcuResult::CCU_SUCCESS;
289	}	183	}
290		184
291	-static bool is_launch_context_valid(const HcommCcuLaunchContextPod& context)
292	-{
293	- const auto& header = context.header;
294	- return header.version == HCOMM_CCU_LAUNCH_ABI_VERSION && header.magicWord == HCOMM_CCU_LAUNCH_CONTEXT_MAGIC_WORD &&
295	- header.size == sizeof(context) && header.reserved == 0 && context.reserved[0] == 0 &&
296	- context.reserved[1] == 0 && context.runtimeStream != 0 && context.threadHandle != 0;
297	-}
298	-
299	CcuResult asccomm_ccu_kernel_launch(	185	CcuResult asccomm_ccu_kernel_launch(
300	- const HcommCcuLaunchContextPod* launch_context, ccu_kernel_handle kernel_handle, const void* task_args,	186	+ ccu_launch_stream stream, ccu_kernel_handle kernel_handle, const void* task_args, uint32_t arg_num)
301	- uint32_t arg_num)
302	{	187	{
303	- const auto start_time = std::chrono::steady_clock::now();	188	+ CCU_CHK_PTR_NULL(stream);
304	-
305	- CCU_CHK_PTR_NULL(launch_context);
306	CHK_PRT_RET(	189	CHK_PRT_RET(
307	kernel_handle == 0, HCCL_ERROR("[%s] failed, kernel handle is empty.", __func__), CcuResult::CCU_E_PARA);	190	kernel_handle == 0, HCCL_ERROR("[%s] failed, kernel handle is empty.", __func__), CcuResult::CCU_E_PARA);
308	CHK_PRT_RET(	191	CHK_PRT_RET(
309	arg_num > 0 && task_args == nullptr,	192	arg_num > 0 && task_args == nullptr,
310	HCCL_ERROR("[%s] failed, taskArgs is nullptr while argNum[%u] > 0.", __func__, arg_num), CcuResult::CCU_E_PTR);	193	HCCL_ERROR("[%s] failed, taskArgs is nullptr while argNum[%u] > 0.", __func__, arg_num), CcuResult::CCU_E_PTR);
311		194
312	- // 跨动态库边界先完整校验 POD，避免按不兼容布局读取 stream、设备号和 profiling 状态。	195	+ // launch 使用调用线程实际所在设备，避免依赖数据面内部设备上下文与调用线程状态一致。
313	- CHK_PRT_RET(	196	+ const int32_t dev_logic_id = asc::get_current_ccu_device_logic_id();
		D D DragonBornHD8418 天前 [提示] 以前设备号是 hcomm 生成 launch context 时写进去的，现在改成"谁调用就取谁当前绑定的设备"：调用线程忘了先 aclrtSetDevice，这里直接报 CCU_E_UNAVAIL；多卡时还可能按另一个设备去查 kernel mgr。建议在头文件接口注释里把这条要求写清楚，使用方一看就知道要先绑设备： `// 注意：launch 按调用线程当前绑定的设备定位 kernel，调用前需已 aclrtSetDevice extern CcuResult asccomm_ccu_kernel_launch( ccu_launch_stream stream, ccu_kernel_handle kernel_handle, const void* task_args, uint32_t arg_num);`
314	- !is_launch_context_valid(*launch_context),	197	+ if (dev_logic_id < 0) {
315	- HCCL_ERROR("[%s] failed, launch context has an incompatible ABI header or runtime stream.", __func__),	198	+ HCCL_ERROR("[%s] failed, current thread has no device set.", __func__);
316	- CcuResult::CCU_E_PARA);	199	+ return CcuResult::CCU_E_UNAVAIL;
317	- HCCL_INFO(	200	+ }
318	- "[asccomm_ccu_kernel_launch] threadHandle[0x%llx] kernelHandle[0x%llx].", launch_context->threadHandle,
319	- kernel_handle);
320	- // runtimeStream 是 hcomm 借出的运行时 stream，且仅在本次同步 launch 返回前有效。
321	- auto stream_ptr = reinterpret_cast<aclrtStream>(launch_context->runtimeStream);
322	-
323	- // launch 使用线程实际所属设备，避免依赖 asc-comm 内部设备上下文与调用线程状态一致。
324	- const int32_t dev_logic_id = launch_context->deviceLogicId;
325	auto& kernel_mgr = asc::ccu_kernel_mgr::get_instance(dev_logic_id);	201	auto& kernel_mgr = asc::ccu_kernel_mgr::get_instance(dev_logic_id);
326	auto* kernel = kernel_mgr.get_kernel(kernel_handle);	202	auto* kernel = kernel_mgr.get_kernel(kernel_handle);
327	CCU_CHK_PTR_NULL(kernel);	203	CCU_CHK_PTR_NULL(kernel);
328		204
		205	+ auto stream_ptr = reinterpret_cast<aclrtStream>(stream);
		206	+
329	CCU_EXCEPTION_HANDLE_BEGIN	207	CCU_EXCEPTION_HANDLE_BEGIN
330	std::vector<asc::ccu_task_param> task_params{};	208	std::vector<asc::ccu_task_param> task_params{};
331	auto ret = kernel->gene_task_params(static_cast<const uint64_t*>(task_args), arg_num, task_params);	209	auto ret = kernel->gene_task_params(static_cast<const uint64_t*>(task_args), arg_num, task_params);
332	CHK_PRT_RET(	210	CHK_PRT_RET(
333	- ret != CcuResult::CCU_SUCCESS,	211	+ ret != CcuResult::CCU_SUCCESS, HCCL_ERROR("[%s] failed, kernelHandle[0x%llx].", __func__, kernel_handle), ret);
334	- HCCL_ERROR(
335	- "[%s] failed, threadHandle[0x%llx] kernelHandle[0x%llx].", __func__, launch_context->threadHandle,
336	- kernel_handle),
337	- ret);
338		212
339	if (task_params.empty()) {	213	if (task_params.empty()) {
340	HCCL_INFO("[%s] passed, ccu params are empty.", __func__);	214	HCCL_INFO("[%s] passed, ccu params are empty.", __func__);
341	return CcuResult::CCU_SUCCESS;	215	return CcuResult::CCU_SUCCESS;
342	}	216	}
343	- const bool is_save_profiling_data = launch_context->profilingL1Enabled != 0 \|\|
344	- launch_context->profilingL0Enabled != 0 \|\| launch_context->profilingCached != 0;
345	-
346	- std::vector<asc::ccu_profiling_info> all_ccu_profiling_info;
347	- CCU_CHK_RET(construct_profiling_info(
348	- kernel, static_cast<const uint64_t*>(task_args), arg_num, all_ccu_profiling_info, is_save_profiling_data));
349	- log_ccu_task_info(task_params, kernel_handle, launch_context->timeoutSec);
350	- std::vector<HcommCcuProfileDetailPod> ccu_detail_info;
351	- CCU_CHK_RET(construct_ccu_detail_info(all_ccu_profiling_info, is_save_profiling_data, ccu_detail_info));
352	for (uint32_t idx = 0; idx < task_params.size(); idx++) {	217	for (uint32_t idx = 0; idx < task_params.size(); idx++) {
353	- uint64_t begin_time = 0;	218	+ CCU_CHK_RET(launch_ccu_tasks(task_params[idx], stream_ptr));
354	- uint64_t end_time = 0;
355	- if (launch_context->getProfilingCycle != nullptr) {
356	- CCU_CHK_RET(launch_context->getProfilingCycle(&begin_time));
357	- }
358	- CCU_CHK_RET(launch_ccu_tasks(task_params[idx], stream_ptr, launch_context->timeoutSec));
359	- if (launch_context->getProfilingCycle != nullptr) {
360	- CCU_CHK_RET(launch_context->getProfilingCycle(&end_time));
361	- }
362	- const HcommCcuTaskProfilePod task_profile = construct_ccu_task_profile(
363	- task_params[idx], kernel_handle, begin_time, end_time, launch_context->isMaster != 0);
364	- if (launch_context->reportTask != nullptr) {
365	- CCU_CHK_RET(launch_context->reportTask(
366	- launch_context->threadHandle, &task_profile, ccu_detail_info.empty() ? nullptr : ccu_detail_info.data(),
367	- static_cast<uint32_t>(ccu_detail_info.size())));
368	- }
369	}	219	}
370	CCU_EXCEPTION_HANDLE_END	220	CCU_EXCEPTION_HANDLE_END
371	- const auto duration =
372	- std::chrono::duration_cast<std::chrono::microseconds>(std::chrono::steady_clock::now() - start_time);
373	- HCCL_INFO("[%s] success, take time [%lld]us.", __func__, duration.count());
374	return CcuResult::CCU_SUCCESS;	221	return CcuResult::CCU_SUCCESS;
375	}	222	}
376		223
	@@ -397,9 +244,9 @@ CcuResult asccomm_ccu_get_mem_token(uint64_t src_va, uint64_t size, uint64_t* to
397	static CcuResult ccu_alloc_var_event_res(	244	static CcuResult ccu_alloc_var_event_res(
398	CcuInsHandle ins_handle, asc::ccu_var_event_type type, uint8_t die_id, uint32_t num, uint64_t& out_handle)	245	CcuInsHandle ins_handle, asc::ccu_var_event_type type, uint8_t die_id, uint32_t num, uint64_t& out_handle)
399	{	246	{
400	- const uint32_t dev_logic_id = asc::get_current_ccu_device_logic_id();	247	+ int32_t dev_logic_id = -1;
401	- auto* ccu_ins = asc::ccu_kernel_registry_mgr::get_instance(dev_logic_id).get(ins_handle);	248	+ asc::ccu_kernel_registry* ccu_ins = nullptr;
402	- CCU_CHK_PTR_NULL(ccu_ins);	249	+ CCU_CHK_RET(ccu_prepare_instance(ins_handle, ccu_ins, dev_logic_id));
		D D DragonBornHD8418 天前 [一般] alloc 路径改成 get_or_create + load 后新开了一个并发窗口：load_register_context 里的 context_loaded_ 判断没有加锁，两个线程（比如一个先调 asccomm_ccu_variable_alloc、另一个同时调 register_start）首次进入时会同时通过判断、并发执行 res_snapshot_->load() 写同一批容器，存在数据竞争崩溃风险（改动前 alloc 只 get 不 load，没有这个窗口）。建议在 registry 内给首次 load 加互斥，或用 std::call_once 保证只加载一次： `std::lock_guard<std::mutex> lock(load_mutex_); // <- ccu_kernel_registry 内加互斥，串行化首次 load if (context_loaded_) { return CcuResult::CCU_SUCCESS; }`
403		250
404	auto* res_pack = ccu_ins->get_res_snapshot();	251	auto* res_pack = ccu_ins->get_res_snapshot();
405	CCU_CHK_PTR_NULL(res_pack);	252	CCU_CHK_PTR_NULL(res_pack);

Msrc/ccu/hcomm/resource/kernel/ccu_kernel_registry.cc+9-4

	@@ -49,11 +49,8 @@ CcuResult ccu_kernel_registry::init(int32_t device_logic_id)
49		49
50	CcuResult ccu_kernel_registry::load_register_context(const HcommCcuInstance& instance)	50	CcuResult ccu_kernel_registry::load_register_context(const HcommCcuInstance& instance)
51	{	51	{
		52	+ std::lock_guard<std::mutex> load_lock(load_mutex_);
52	CCU_EXCEPTION_HANDLE_BEGIN	53	CCU_EXCEPTION_HANDLE_BEGIN
53	- if (!res_snapshot_) {
54	- res_snapshot_.reset(new (std::nothrow) asc_ccu_res_snapshot());
55	- CCU_CHK_PTR_NULL(res_snapshot_);
56	- }
57	// 当前线程设备必须与创建 registry 时一致，否则说明上下文被用在了错误的设备上	54	// 当前线程设备必须与创建 registry 时一致，否则说明上下文被用在了错误的设备上
58	// （Instance 不再携带设备号，定档口径：即时取 aclrtGetDevice）。	55	// （Instance 不再携带设备号，定档口径：即时取 aclrtGetDevice）。
59	const int32_t current_device = asc::get_current_ccu_device_logic_id();	56	const int32_t current_device = asc::get_current_ccu_device_logic_id();
	@@ -63,7 +60,15 @@ CcuResult ccu_kernel_registry::load_register_context(const HcommCcuInstance& ins
63	current_device, dev_logic_id_);	60	current_device, dev_logic_id_);
64	return CcuResult::CCU_E_PARA;	61	return CcuResult::CCU_E_PARA;
65	}	62	}
		63	+ if (context_loaded_) {
		64	+ return CcuResult::CCU_SUCCESS;
		65	+ }
		66	+ if (!res_snapshot_) {
		67	+ res_snapshot_.reset(new (std::nothrow) asc_ccu_res_snapshot());
		68	+ CCU_CHK_PTR_NULL(res_snapshot_);
		69	+ }
66	CCU_CHK_RET(res_snapshot_->load(instance));	70	CCU_CHK_RET(res_snapshot_->load(instance));
		71	+ context_loaded_ = true;
67	CCU_EXCEPTION_HANDLE_END	72	CCU_EXCEPTION_HANDLE_END
68	return CcuResult::CCU_SUCCESS;	73	return CcuResult::CCU_SUCCESS;
69	}	74	}

Msrc/ccu/hcomm/resource/kernel/ccu_kernel_registry.h+5-0

	@@ -12,6 +12,7 @@
12	#define CCU_KERNEL_REGISTRY_H	12	#define CCU_KERNEL_REGISTRY_H
13		13
14	#include <memory>	14	#include <memory>
		15	+#include <mutex>
15	#include <vector>	16	#include <vector>
16		17
17	#include "ccu/hcomm/ccu_api_types.h"	18	#include "ccu/hcomm/ccu_api_types.h"
	@@ -54,6 +55,10 @@ private:
54	// CcuVarEventResMgr 按 handle 记账，析构时需以此调用 ReleaseByInstance。	55	// CcuVarEventResMgr 按 handle 记账，析构时需以此调用 ReleaseByInstance。
55	CcuInsHandle ins_handle_{};	56	CcuInsHandle ins_handle_{};
56	std::unique_ptr<asc_ccu_res_snapshot> res_snapshot_{};	57	std::unique_ptr<asc_ccu_res_snapshot> res_snapshot_{};
		58	+ // 资源接口可能在首次 kernel 注册前加载 Instance；后续 <<<>>> 注册轮复用该快照，避免重复加载。
		59	+ // 该锁仅串行化首次快照加载；同一实例的注册状态和资源申请状态仍要求调用方串行使用。
		60	+ std::mutex load_mutex_{};
		61	+ bool context_loaded_{false};
57	std::vector<ccu_kernel_handle> kernel_handles_{};	62	std::vector<ccu_kernel_handle> kernel_handles_{};
58	std::vector<ccu_kernel_handle> untranslated_kernel_handles_{};	63	std::vector<ccu_kernel_handle> untranslated_kernel_handles_{};
59	};	64	};

Mtests/ut/CMakeLists.txt+10-0

	@@ -280,9 +280,19 @@ target_compile_options(asccomm_ut_ccu_host_launch PRIVATE
280	-Wall	280	-Wall
281	-Werror	281	-Werror
282	)	282	)
		283	+
		284	+add_library(asccomm_ccu_launch_stub STATIC
		285	+ ${ASCCOMM_UT_ROOT}/ccu/stub/ccu_launch_stub.cc
		286	+)
		287	+target_include_directories(asccomm_ccu_launch_stub PRIVATE
		288	+ ${ASCCOMM_ROOT}/include
		289	+ ${ASCEND_CANN_ARCH_DIR}/include
		290	+)
		291	+
283	target_link_libraries(asccomm_ut_ccu_host_launch PRIVATE	292	target_link_libraries(asccomm_ut_ccu_host_launch PRIVATE
284	intf_llt_pub_basic	293	intf_llt_pub_basic
285	asccomm_ccu_host_launch	294	asccomm_ccu_host_launch
		295	+ asccomm_ccu_launch_stub
286	)	296	)
287		297
288	run_llt_test(	298	run_llt_test(

Atests/ut/ccu/stub/ccu_launch_stub.cc+124-0

@@ -0,0 +1,124 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include "ccu_launch_stub.h"
	12	+
	13	+namespace {
	14	+
	15	+CcuResult g_registerStartResult = CcuResult::CCU_SUCCESS;
	16	+CcuResult g_registerKernelResult = CcuResult::CCU_SUCCESS;
	17	+CcuResult g_registerEndResult = CcuResult::CCU_SUCCESS;
	18	+CcuResult g_getTaskArgsNumResult = CcuResult::CCU_SUCCESS;
	19	+CcuResult g_kernelLaunchResult = CcuResult::CCU_SUCCESS;
	20	+uint32_t g_taskArgsNum = 0;
	21	+CcuLaunchStubObservations g_observations{};
	22	+
	23	+} // namespace
	24	+
	25	+void ResetCcuLaunchStub()
	26	+{
	27	+ g_registerStartResult = CcuResult::CCU_SUCCESS;
	28	+ g_registerKernelResult = CcuResult::CCU_SUCCESS;
	29	+ g_registerEndResult = CcuResult::CCU_SUCCESS;
	30	+ g_getTaskArgsNumResult = CcuResult::CCU_SUCCESS;
	31	+ g_kernelLaunchResult = CcuResult::CCU_SUCCESS;
	32	+ g_taskArgsNum = 0;
	33	+ g_observations = {};
	34	+}
	35	+
	36	+void SetCcuLaunchStubResults(
	37	+ CcuResult registerStart, CcuResult registerKernel, CcuResult registerEnd, CcuResult getTaskArgsNum,
	38	+ CcuResult kernelLaunch)
	39	+{
	40	+ g_registerStartResult = registerStart;
	41	+ g_registerKernelResult = registerKernel;
	42	+ g_registerEndResult = registerEnd;
	43	+ g_getTaskArgsNumResult = getTaskArgsNum;
	44	+ g_kernelLaunchResult = kernelLaunch;
	45	+}
	46	+
	47	+void SetCcuLaunchStubTaskArgsNum(uint32_t taskArgsNum) { g_taskArgsNum = taskArgsNum; }
	48	+
	49	+CcuLaunchStubObservations GetCcuLaunchStubObservations() { return g_observations; }
	50	+
	51	+extern "C" CcuResult asccomm_ccu_kernel_register_start(CcuInsHandle insHandle)
	52	+{
	53	+ ++g_observations.calls.registerStart;
	54	+ g_observations.lastInsHandle = insHandle;
	55	+ return g_registerStartResult;
	56	+}
	57	+
	58	+extern "C" CcuResult asccomm_ccu_kernel_register(
	59	+ CcuInsHandle insHandle, uint32_t dieId, const char* kernelFuncName, const void* kernelFunc, const void** kernelArgs,
	60	+ uint32_t argNum, ccu_kernel_handle* kernelHandle)
	61	+{
	62	+ ++g_observations.calls.registerKernel;
	63	+ g_observations.lastInsHandle = insHandle;
	64	+ g_observations.lastDieId = dieId;
	65	+ g_observations.lastKernelName = kernelFuncName;
	66	+ g_observations.lastKernelFunc = kernelFunc;
	67	+ if (kernelArgs != nullptr && argNum > 0) {
	68	+ g_observations.lastKernelArg = kernelArgs[0];
	69	+ }
	70	+ if (kernelHandle == nullptr) {
	71	+ return CcuResult::CCU_E_PTR;
	72	+ }
	73	+ if (g_registerKernelResult != CcuResult::CCU_SUCCESS) {
	74	+ return g_registerKernelResult;
	75	+ }
	76	+
	77	+ // The wrapper registers a trampoline returning CcuResult. Execute it here to cover the
	78	+ // void-kernel adaptation without depending on the real kernel representation stack.
	79	+ if (argNum != 1U \|\| kernelArgs == nullptr) {
	80	+ return CcuResult::CCU_E_PARA;
	81	+ }
	82	+ using CcuKernelFunc = CcuResult (*)(ccu_kernel_arg);
	83	+ auto kernel = reinterpret_cast<CcuKernelFunc>(const_cast<void*>(kernelFunc));
	84	+ CcuResult ret = kernel(const_cast<void*>(kernelArgs[0]));
	85	+ if (ret != CcuResult::CCU_SUCCESS) {
	86	+ return ret;
	87	+ }
	88	+ *kernelHandle = 0x5100U;
	89	+ g_observations.lastKernelHandle = *kernelHandle;
	90	+ return CcuResult::CCU_SUCCESS;
	91	+}
	92	+
	93	+extern "C" CcuResult asccomm_ccu_kernel_register_end(CcuInsHandle insHandle)
	94	+{
	95	+ ++g_observations.calls.registerEnd;
	96	+ g_observations.lastInsHandle = insHandle;
	97	+ return g_registerEndResult;
	98	+}
	99	+
	100	+extern "C" CcuResult asccomm_ccu_get_task_args_num(ccu_kernel_handle kernelHandle, uint32_t* taskArgsNum)
	101	+{
	102	+ ++g_observations.calls.getTaskArgsNum;
	103	+ g_observations.lastKernelHandle = kernelHandle;
	104	+ if (taskArgsNum == nullptr) {
	105	+ return CcuResult::CCU_E_PTR;
	106	+ }
	107	+ if (g_getTaskArgsNumResult != CcuResult::CCU_SUCCESS) {
	108	+ return g_getTaskArgsNumResult;
	109	+ }
	110	+ *taskArgsNum = g_taskArgsNum;
	111	+ g_observations.lastTaskArgsNum = *taskArgsNum;
	112	+ return CcuResult::CCU_SUCCESS;
	113	+}
	114	+
	115	+extern "C" CcuResult asccomm_ccu_kernel_launch(
	116	+ ccu_launch_stream stream, ccu_kernel_handle kernelHandle, const void* taskArgs, uint32_t argNum)
	117	+{
	118	+ ++g_observations.calls.kernelLaunch;
	119	+ g_observations.lastStream = stream;
	120	+ g_observations.lastKernelHandle = kernelHandle;
	121	+ g_observations.lastTaskArgs = taskArgs;
	122	+ g_observations.lastLaunchArgNum = argNum;
	123	+ return g_kernelLaunchResult;
	124	+}

Atests/ut/ccu/stub/ccu_launch_stub.h+45-0

@@ -0,0 +1,45 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software: you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef CCU_LAUNCH_STUB_H
	12	+#define CCU_LAUNCH_STUB_H
	13	+
	14	+#include "ccu/hcomm/ccu_launch.h"
	15	+
	16	+struct CcuLaunchStubCalls {
	17	+ uint32_t registerStart;
	18	+ uint32_t registerKernel;
	19	+ uint32_t registerEnd;
	20	+ uint32_t getTaskArgsNum;
	21	+ uint32_t kernelLaunch;
	22	+};
	23	+
	24	+struct CcuLaunchStubObservations {
	25	+ CcuLaunchStubCalls calls{};
	26	+ CcuInsHandle lastInsHandle{};
	27	+ uint32_t lastDieId{};
	28	+ const char* lastKernelName{};
	29	+ const void* lastKernelFunc{};
	30	+ const void* lastKernelArg{};
	31	+ ccu_kernel_handle lastKernelHandle{};
	32	+ uint32_t lastTaskArgsNum{};
	33	+ ccu_launch_stream lastStream{};
	34	+ const void* lastTaskArgs{};
	35	+ uint32_t lastLaunchArgNum{};
	36	+};
	37	+
	38	+void ResetCcuLaunchStub();
	39	+void SetCcuLaunchStubResults(
	40	+ CcuResult registerStart, CcuResult registerKernel, CcuResult registerEnd, CcuResult getTaskArgsNum,
	41	+ CcuResult kernelLaunch);
	42	+void SetCcuLaunchStubTaskArgsNum(uint32_t taskArgsNum);
	43	+CcuLaunchStubObservations GetCcuLaunchStubObservations();
	44	+
	45	+#endif // CCU_LAUNCH_STUB_H

Mtests/ut/ccu/test_ccu_host_launch.cpp+189-199

 #include <gtest/gtest.h>
--#include "acl/acl_rt.h"
++#include <vector>
++
 #include "ccu/ccu_host_launch.h"
++#include "stub/ccu_launch_stub.h"
 namespace {
--uint32_t g_registerStartCallCount = 0;
--CcuKernelHandle g_nextKernelHandle = 1;
--uint32_t g_threadAllocCallCount = 0;
--uint32_t g_threadFreeCallCount = 0;
--uint32_t g_kernelLaunchCallCount = 0;
--uintptr_t g_nextStream = 0x22;
--CommEngine g_observedThreadEngine = COMM_ENGINE_RESERVED;
--aclrtStream g_observedThreadStream = nullptr;
--ThreadHandle g_nextThreadHandle = 0x100;
--ThreadHandle g_observedLaunchThread = 0;
--CcuKernelHandle g_observedLaunchKernel = 0;
--const void* g_observedLaunchArgs = nullptr;
--uint32_t g_observedLaunchArgNum = 0;
--uint32_t g_taskArgsNum = 3;
--const char* g_observedRegisterKernelName = nullptr;
--CcuResult g_kernelLaunchResult = CCU_SUCCESS;
--HcommResult g_threadAllocResult = HCCL_SUCCESS;
--HcommResult g_threadFreeResult = HCCL_SUCCESS;
--
--void DummyKernel(void*) {}
--
 constexpr const char* DUMMY_KERNEL_NAME = "DummyKernel";
++uint32_t g_dummyKernelCallCount = 0;
++const void* g_dummyKernelObservedArg = nullptr;
++
++void DummyKernel(void* arg)
++{
++    ++g_dummyKernelCallCount;
++    g_dummyKernelObservedArg = arg;
++}
 asccomm_launch_kernel_cfg MakeLaunchCfg()
 {
     asccomm_launch_kernel_cfg cfg{};
     cfg.ccu_schd = {1, 0, 0x01, 0};
--    cfg.ccu_ins = 0x11;
++    cfg.ccu_ins = CcuInsHandle{};
--    cfg.stream = reinterpret_cast<aclrtStream>(g_nextStream++);
++    cfg.stream = reinterpret_cast<aclrtStream>(0x22);
     cfg.attrs = nullptr;
     return cfg;
 }
 } // namespace
--extern "C" CcuResult HcommCcuKernelRegisterStart(CcuInsHandle)
--{
--    ++g_registerStartCallCount;
--    return CCU_SUCCESS;
--}
--
--extern "C" CcuResult HcommCcuKernelRegister(
--    CcuInsHandle, uint32_t, const char* kernelName, const void*, const void**, uint32_t, CcuKernelHandle* kernelHandle)
--{
--    g_observedRegisterKernelName = kernelName;
--    if (kernelHandle == nullptr) {
--        return CCU_E_PTR;
--    }
--    *kernelHandle = g_nextKernelHandle++;
--    return CCU_SUCCESS;
--}
--
--extern "C" CcuResult HcommCcuKernelRegisterEnd(CcuInsHandle) { return CCU_SUCCESS; }
--
--extern "C" CcuResult HcommCcuGetTaskArgsNum(CcuKernelHandle, uint32_t* taskArgsNum)
--{
--    if (taskArgsNum == nullptr) {
--        return CCU_E_PTR;
--    }
--    *taskArgsNum = g_taskArgsNum;
--    return CCU_SUCCESS;
--}
--
--extern "C" HcommResult HcommThreadAllocWithStream(CommEngine engine, aclrtStream stream, uint32_t, ThreadHandle* thread)
--{
--    ++g_threadAllocCallCount;
--    g_observedThreadEngine = engine;
--    g_observedThreadStream = stream;
--    if (g_threadAllocResult == HCCL_SUCCESS && thread != nullptr) {
--        *thread = g_nextThreadHandle++;
--    }
--    return g_threadAllocResult;
--}
--
--extern "C" HcommResult HcommThreadFree(const ThreadHandle*, uint32_t)
--{
--    ++g_threadFreeCallCount;
--    return g_threadFreeResult;
--}
--
--extern "C" CcuResult HcommCcuKernelLaunch(
--    ThreadHandle thread, CcuKernelHandle kernel, const void* launchArgs, uint32_t argNum)
--{
--    ++g_kernelLaunchCallCount;
--    g_observedLaunchThread = thread;
--    g_observedLaunchKernel = kernel;
--    g_observedLaunchArgs = launchArgs;
--    g_observedLaunchArgNum = argNum;
--    return g_kernelLaunchResult;
--}
--
 class TestCcuHostLaunch : public testing::Test {
 protected:
     void SetUp() override
     {
--        g_registerStartCallCount = 0;
++        ResetCcuLaunchStub();
--        g_threadAllocCallCount = 0;
++        SetCcuLaunchStubTaskArgsNum(3U);
--        g_threadFreeCallCount = 0;
++        g_dummyKernelCallCount = 0;
--        g_kernelLaunchCallCount = 0;
++        g_dummyKernelObservedArg = nullptr;
--        g_observedThreadEngine = COMM_ENGINE_RESERVED;
--        g_observedThreadStream = nullptr;
--        g_observedLaunchThread = 0;
--        g_observedLaunchKernel = 0;
--        g_observedLaunchArgs = nullptr;
--        g_observedLaunchArgNum = 0;
--        g_observedRegisterKernelName = nullptr;
--        g_taskArgsNum = 3;
--        g_kernelLaunchResult = CCU_SUCCESS;
--        g_threadAllocResult = HCCL_SUCCESS;
--        g_threadFreeResult = HCCL_SUCCESS;
     }
 };
         CCU_E_PARA);
 }
--TEST_F(TestCcuHostLaunch, LaunchUsesAllocatedThreadAndHcommKernelLaunch)
--{
--    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
--    uint64_t args[] = {1, 2, 3};
--
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
--        CCU_SUCCESS);
--    EXPECT_EQ(g_threadAllocCallCount, 1U);
--    EXPECT_EQ(g_observedThreadEngine, COMM_ENGINE_CCU);
--    EXPECT_EQ(g_observedThreadStream, cfg.stream);
--    EXPECT_EQ(g_kernelLaunchCallCount, 1U);
--    EXPECT_NE(g_observedLaunchThread, 0U);
--    EXPECT_NE(g_observedLaunchKernel, 0U);
--    EXPECT_EQ(g_observedLaunchArgs, args);
--    EXPECT_EQ(g_observedLaunchArgNum, 3U);
--    EXPECT_EQ(g_observedRegisterKernelName, DUMMY_KERNEL_NAME);
--    EXPECT_EQ(g_threadFreeCallCount, 0U);
--}
--
--TEST_F(TestCcuHostLaunch, NullKernelNameIsForwardedToRegister)
--{
--    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
--    uint64_t args[] = {1, 2, 3};
--
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), nullptr, &cfg, args), CCU_SUCCESS);
--    EXPECT_EQ(g_observedRegisterKernelName, nullptr);
--    EXPECT_EQ(g_kernelLaunchCallCount, 1U);
--}
--
--TEST_F(TestCcuHostLaunch, TooManyTaskArgsReturnsNotSupportBeforeThreadAllocation)
--{
--    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
--    uint64_t args[] = {1};
--    g_taskArgsNum = 14;
--
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
--        CCU_E_NOT_SUPPORT);
--    EXPECT_EQ(g_threadAllocCallCount, 0U);
--    EXPECT_EQ(g_kernelLaunchCallCount, 0U);
--}
--
--TEST_F(TestCcuHostLaunch, ThreadAllocationFailureStopsLaunch)
--{
--    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
--    uint64_t args[] = {1, 2, 3};
--    g_threadAllocResult = HCCL_E_UNAVAIL;
--
--    // MakeLaunchCfg uses a fresh stream, so this test does not reuse a cached handle.
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
--        CCU_E_UNAVAIL);
--    EXPECT_EQ(g_kernelLaunchCallCount, 0U);
--    EXPECT_EQ(g_threadFreeCallCount, 0U);
--}
--
--TEST_F(TestCcuHostLaunch, KernelLaunchFailureKeepsThreadAlive)
--{
--    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
--    uint64_t args[] = {1, 2, 3};
--    g_kernelLaunchResult = CCU_E_RUNTIME;
--
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
--        CCU_E_RUNTIME);
--    EXPECT_EQ(g_kernelLaunchCallCount, 1U);
--    EXPECT_EQ(g_threadFreeCallCount, 0U);
--}
--
--TEST_F(TestCcuHostLaunch, RepeatedLaunchesAllocateAndFreeIndependentThreads)
--{
--    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
--    uint64_t args[] = {1, 2, 3};
--
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
--        CCU_SUCCESS);
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
--        CCU_SUCCESS);
--    EXPECT_EQ(g_threadAllocCallCount, 1U);
--    EXPECT_EQ(g_kernelLaunchCallCount, 2U);
--    EXPECT_EQ(g_threadFreeCallCount, 0U);
--}
--
--TEST_F(TestCcuHostLaunch, DifferentStreamsUseIndependentThreads)
--{
--    asccomm_launch_kernel_cfg firstCfg = MakeLaunchCfg();
--    asccomm_launch_kernel_cfg secondCfg = MakeLaunchCfg();
--    uint64_t args[] = {1, 2, 3};
--
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &firstCfg, args),
--        CCU_SUCCESS);
--    EXPECT_EQ(
--        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &secondCfg, args),
--        CCU_SUCCESS);
--    EXPECT_EQ(g_threadAllocCallCount, 2U);
--    EXPECT_EQ(g_kernelLaunchCallCount, 2U);
--    EXPECT_EQ(g_threadFreeCallCount, 0U);
--}
--
 TEST_F(TestCcuHostLaunch, LaunchHashTagReturnsStableNonZeroValue)
 {
     const uint64_t tag = asccomm_ccu_get_launch_hash_tag("CcuAllGatherMesh1DMem2MemKernel");
     EXPECT_EQ(tag, HcommCcuGetLaunchHashTag("CcuAllGatherMesh1DMem2MemKernel"));
 }
++TEST_F(TestCcuHostLaunch, LaunchSuccessForwardsRegisterAndLaunchParameters)
++{
++    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++    cfg.ccu_schd.binary_cache_tag = 0x5101U;
++    uint64_t args[] = {1, 2, 3};
++
++    EXPECT_EQ(
++        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++        CCU_SUCCESS);
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.calls.registerStart, 1U);
++    EXPECT_EQ(observed.calls.registerKernel, 1U);
++    EXPECT_EQ(observed.calls.registerEnd, 1U);
++    EXPECT_EQ(observed.calls.getTaskArgsNum, 1U);
++    EXPECT_EQ(observed.calls.kernelLaunch, 1U);
++    EXPECT_EQ(observed.lastInsHandle.ccuInsKey, cfg.ccu_ins.ccuInsKey);
++    EXPECT_EQ(observed.lastInsHandle.ccuInsPtr, cfg.ccu_ins.ccuInsPtr);
++    EXPECT_EQ(observed.lastDieId, 0U);
++    EXPECT_EQ(observed.lastKernelName, DUMMY_KERNEL_NAME);
++    EXPECT_NE(observed.lastKernelHandle, 0U);
++    EXPECT_EQ(observed.lastTaskArgsNum, 3U);
++    EXPECT_EQ(observed.lastStream, cfg.stream);
++    EXPECT_EQ(observed.lastTaskArgs, args);
++    EXPECT_EQ(observed.lastLaunchArgNum, 3U);
++    EXPECT_EQ(g_dummyKernelCallCount, 1U);
++    EXPECT_EQ(g_dummyKernelObservedArg, args);
++}
++
++TEST_F(TestCcuHostLaunch, NullKernelNameIsForwardedToRegister)
++{
++    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++    uint64_t args[] = {1, 2, 3};
++
++    EXPECT_EQ(
++        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), nullptr, &cfg, args), CCU_SUCCESS);
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.lastKernelName, nullptr);
++    EXPECT_EQ(observed.calls.kernelLaunch, 1U);
++}
++
++TEST_F(TestCcuHostLaunch, RegisterFailureClosesRegisterRound)
++{
++    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++    uint64_t args[] = {1, 2, 3};
++    SetCcuLaunchStubResults(
++        CcuResult::CCU_SUCCESS, CcuResult::CCU_E_RUNTIME, CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS,
++        CcuResult::CCU_SUCCESS);
++
++    EXPECT_EQ(
++        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++        CCU_E_RUNTIME);
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.calls.registerStart, 1U);
++    EXPECT_EQ(observed.calls.registerKernel, 1U);
++    EXPECT_EQ(observed.calls.registerEnd, 1U);
++    EXPECT_EQ(observed.calls.getTaskArgsNum, 0U);
++    EXPECT_EQ(observed.calls.kernelLaunch, 0U);
++}
++
++TEST_F(TestCcuHostLaunch, LaunchFailureIsPropagatedAndCachedHandleIsReused)
++{
++    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++    cfg.ccu_schd.binary_cache_tag = 0x5301U;
++    uint64_t args[] = {1, 2, 3};
++    SetCcuLaunchStubResults(
++        CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS,
++        CcuResult::CCU_E_RUNTIME);
++
++    EXPECT_EQ(
++        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++        CCU_E_RUNTIME);
++
++    SetCcuLaunchStubResults(
++        CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS, CcuResult::CCU_SUCCESS,
++        CcuResult::CCU_SUCCESS);
++    EXPECT_EQ(
++        asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++        CCU_SUCCESS);
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.calls.registerStart, 1U);
++    EXPECT_EQ(observed.calls.registerKernel, 1U);
++    EXPECT_EQ(observed.calls.registerEnd, 1U);
++    EXPECT_EQ(observed.calls.getTaskArgsNum, 2U);
++    EXPECT_EQ(observed.calls.kernelLaunch, 2U);
++}
++
++TEST_F(TestCcuHostLaunch, SameCacheTagRegistersOnceAndLaunchesEachTime)
++{
++    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++    cfg.ccu_schd.binary_cache_tag = 0x5401U;
++    uint64_t args[] = {1, 2, 3};
++
++    for (uint32_t idx = 0; idx < 2; idx++) {
++        EXPECT_EQ(
++            asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++            CCU_SUCCESS);
++    }
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.calls.registerStart, 1U);
++    EXPECT_EQ(observed.calls.registerKernel, 1U);
++    EXPECT_EQ(observed.calls.registerEnd, 1U);
++    EXPECT_EQ(observed.calls.getTaskArgsNum, 2U);
++    EXPECT_EQ(observed.calls.kernelLaunch, 2U);
++}
++
++TEST_F(TestCcuHostLaunch, ZeroCacheTagRegistersEachLaunch)
++{
++    asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++    cfg.ccu_schd.binary_cache_tag = 0U;
++    uint64_t args[] = {1, 2, 3};
++
++    for (uint32_t idx = 0; idx < 2; idx++) {
++        EXPECT_EQ(
++            asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++            CCU_SUCCESS);
++    }
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.calls.registerStart, 2U);
++    EXPECT_EQ(observed.calls.registerKernel, 2U);
++    EXPECT_EQ(observed.calls.registerEnd, 2U);
++    EXPECT_EQ(observed.calls.getTaskArgsNum, 2U);
++    EXPECT_EQ(observed.calls.kernelLaunch, 2U);
++}
++
++TEST_F(TestCcuHostLaunch, DifferentCacheTagsRegisterIndependently)
++{
++    uint64_t args[] = {1, 2, 3};
++    for (uint64_t cacheTag = 0x5501U; cacheTag <= 0x5502U; cacheTag++) {
++        asccomm_launch_kernel_cfg cfg = MakeLaunchCfg();
++        cfg.ccu_schd.binary_cache_tag = cacheTag;
++        EXPECT_EQ(
++            asccomm_ccu_host_kernel_launch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &cfg, args),
++            CCU_SUCCESS);
++    }
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.calls.registerStart, 2U);
++    EXPECT_EQ(observed.calls.registerKernel, 2U);
++    EXPECT_EQ(observed.calls.registerEnd, 2U);
++    EXPECT_EQ(observed.calls.getTaskArgsNum, 2U);
++    EXPECT_EQ(observed.calls.kernelLaunch, 2U);
++}
++
++TEST_F(TestCcuHostLaunch, HcompatCfgForwardsThroughNewLaunchChain)
++{
++    HcommLaunchKernelCfg hcompatCfg{};
++    hcompatCfg.ccuSchd.numBlocks = 1;
++    hcompatCfg.ccuSchd.reserved = 0;
++    hcompatCfg.ccuSchd.phyDieMask = 0x02;
++    hcompatCfg.ccuSchd.binaryCacheTag = 0x5601U;
++    hcompatCfg.ccuIns = CcuInsHandle{};
++    hcompatCfg.stream = reinterpret_cast<aclrtStream>(0x33);
++    hcompatCfg.attrs = nullptr;
++    uint64_t args[] = {1, 2, 3};
++
++    EXPECT_EQ(
++        HcommCcuHostKernelLaunch(reinterpret_cast<const void*>(DummyKernel), DUMMY_KERNEL_NAME, &hcompatCfg, args),
++        CCU_SUCCESS);
++
++    const CcuLaunchStubObservations observed = GetCcuLaunchStubObservations();
++    EXPECT_EQ(observed.lastDieId, 1U);
++    EXPECT_EQ(observed.lastStream, hcompatCfg.stream);
++    EXPECT_EQ(observed.lastTaskArgs, args);
++    EXPECT_EQ(observed.lastLaunchArgNum, 3U);
++}
++
 int main(int argc, char** argv)
 {
     testing::InitGoogleTest(&argc, argv);