BatchWrite
产品支持情况
|
产品 |
是否支持 |
|---|---|
|
|
x |
|
|
√ |
|
|
x |
|
|
x |
|
|
x |
|
|
x |
|
|
x |
功能说明
集合通信BatchWrite的任务下发接口，返回该任务的标识handleId给用户。
BatchWrite实现了一种不同昇腾AI Server（通常是8卡或16卡的昇腾NPU设备组成的服务器形态的统称）间批量的点对点通信，即多卡环境下，不同AI Server间的卡直接传输数据的通信模式。本接口支持将本卡的多份数据同时发送到不同AI Server的多张卡上。
函数原型
1 2 |
template <bool commit = false> __aicore__ inline HcclHandle BatchWrite(GM_ADDR batchWriteInfo, uint32_t itemNum) |
参数说明
|
参数名 |
输入/输出 |
描述 |
|---|---|---|
|
commit |
输入 |
bool类型。参数取值如下：
|
|
参数名 |
输入/输出 |
描述 |
||
|---|---|---|---|---|
|
batchWriteInfo |
输入 |
通信任务信息的Global Memory地址。一组通信数据的相关信息必须按如下格式保存，在执行通信任务时，可以同时指定多组通信任务信息，执行通信任务时批量发送数据。
|
||
|
itemNum |
输入 |
批量任务的个数。该参数取值必须与batchWriteInfo中通信任务信息的组数一致。 |
返回值说明
返回该任务的标识handleId，handleId大于等于0。调用失败时，返回 -1。
约束说明
- 调用本接口前确保已调用过Init和SetCcTiling接口。
- 若Hccl对象的config模板参数未指定下发通信任务的核，该接口只能在AI Cube核或者AI Vector核两者之一上调用。若Hccl对象的config模板参数中指定了下发通信任务的核，则该接口可以在AI Cube核和AI Vector核上同时调用，接口内部会根据指定的核的类型，只在AI Cube核、AI Vector核二者之一下发该通信任务。
- 一个通信域内，所有Prepare接口和InterHcclGroupSync接口的总调用次数不能超过63。
- 当前接口仅支持不同AI Server间的通信，同时通信任务信息中指定的目的卡号不能是本卡号。
- 通信任务信息写入batchWriteInfo时，必须通过调用DataCacheCleanAndInvalid接口，保证预期的数据成功刷新到Global Memory上。
调用示例
假设本卡要将不同的数据分别发送到其它AI Server的2卡、3卡的指定位置，通过调用一次Batchwrite接口，实现批量点对点通信。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 |
struct BatchWriteItem { uint64_t localBuf; // 本端发送数据的window地址 uint64_t remoteBuf; // 对端接收数据的window地址 uint64_t count; // 发送的数据个数 uint32_t dataType; // 发送的数据类型 uint32_t remoteRankId; // 发送数据的目的卡号 }; // 按接口的约定定义格式 extern "C" __global__ __aicore__ void BatchWrite_custom(GM_ADDR inputGM, GM_ADDR workspace, GM_ADDR tilingGM) { GM_ADDR userWS = GetUserWorkspace(workspace); if (userWS == nullptr) { return; } REGISTER_TILING_DEFAULT(BatchWriteCustomTilingData); // BatchWriteCustomTilingData为对应算子头文件定义的结构体 auto tiling = (__gm__ BatchWriteCustomTilingData *)tilingGM; GM_ADDR contextGM = AscendC::GetHcclContext<0>(); __gm__ void *mc2InitTiling = (__gm__ void *)(&tiling->mc2InitTiling); __gm__ void *batchWriteTiling = (__gm__ void *)(&(tiling->mc2CcTiling)); if constexpr (g_coreType == AscendC::AIV) { Hccl hccl; hccl.Init(contextGM, mc2InitTiling); hccl.SetCcTiling(batchWriteTiling); __gm__ BatchWriteItem *sendInfo = reinterpret_cast<__gm__ BatchWriteItem *>(workspace); // 需要提前将待发送的数据从inputGM搬运到localBuf所填的window地址上 sendInfo->localBuf = hccl.GetWindowsOutAddr(hccl.GetRankId()); // 对端的接收地址也要是window地址，接收端需要考虑是否搬运到输出或者workspace上 sendInfo->remoteBuf = hccl.GetWindowsInAddr(2U); sendInfo->count = 16U; sendInfo->dataType = HcclDataType::HCCL_DATA_TYPE_FP16; sendInfo->remoteRankId = 2U; // 可以组装多个通信任务，实现批量发送 (sendInfo + 1)->localBuf = hccl.GetWindowsOutAddr(hccl.GetRankId()); (sendInfo + 1)->remoteBuf = hccl.GetWindowsInAddr(3U); (sendInfo + 1)->count = 32U; (sendInfo + 1)->dataType = HcclDataType::HCCL_DATA_TYPE_BFP16; (sendInfo + 1)->remoteRankId = 3U; // 确保cache中的数据已刷新到GM地址上 GlobalTensor<int64_t> tempTensor; tempTensor.SetGlobalBuffer((__gm__ int64_t *)sendInfo); DataCacheCleanAndInvalid<int64_t, CacheLine::SINGLE_CACHE_LINE, DcciDst::CACHELINE_OUT>(tempTensor); auto handleId = hccl.BatchWrite<true>(sendInfo, 2U); // wait仅表示本端发送完毕，对端是否接收到数据需要在对端判断 hccl.Wait(handleId); AscendC::SyncAll(); hccl.Finalize(); } } |
当通信数据量较大时，可以在Tiling流程中调用SetAicpuBlockDim接口来设置AI CPU的核数。算子内部将自动在多个AI CPU核中选择最优的核进行通信，以实现更优的性能。建议将可调度的AI CPU核数设置为5。
static ge::graphStatus BatchWriteTilingFunc(gert::TilingContext* context)
{
// 省略无关代码
auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
const auto aicCoreNum = ascendcPlatform.GetCoreNumAic();
auto coreNum = use_aiv ? aicCoreNum * 2 : aicCoreNum;
context->SetAicpuBlockDim(5U);
context->SetBlockDim(coreNum);
context->SetTilingKey(1000);
// 省略无关代码
SdmaBatchWriteCustomTilingData *tiling = context->GetTilingData<SdmaBatchWriteCustomTilingData>();
AscendC::Mc2CcTilingConfig mc2CcTilingConfig(groupName, 18, "BatchWrite=level0:fullmesh", 0);
mc2CcTilingConfig.GetTiling(tiling->mc2InitTiling);
mc2CcTilingConfig.GetTiling(tiling->mc2CcTiling);
return ge::GRAPH_SUCCESS;
}
父主题： Hccl