cann/runtime 多设备编程

多设备编程

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime

多设备编程总述

多设备编程使用户能够利用多个NPU（Neural-Network Processing Unit）的综合性能和内存等资源，实现超越单个NPU的性能水平。通常，每个物理NPU在Runtime编程界面中被抽象为一个Device，任务下发时需要有Device中的Context进行支撑，因此多设备编程需要管理多个Device及其相应的Context。

一些常见的多NPU编程方法包括：

• 单个主机线程驱动多个NPU。
• 多个主机线程，每个线程驱动自己的NPU。
• 多个单线程主机进程，每个进程驱动自己的NPU。
• 多个主机进程，每个进程包含多个线程，每个线程驱动自己的NPU。

多Device选择

一个Host搭配多Device的场景下，用户可以在Host侧应用程序中通过aclrtGetDeviceCount接口来获取当前Host上搭配的Device数量，Device按照0、1、2、... 的顺序排布。

以下为获取Device信息的代码示例，不可以直接拷贝编译运行，仅供参考：

// 获取Device数量及其对应的属性信息 uint32_t deviceCount; aclrtGetDeviceCount(&deviceCount); uint32_t deviceId; for (deviceId = 0; deviceId < deviceCount; ++deviceId) { // 按需查询设备属性信息 aclrtDevAttr attr = ACL_DEV_ATTR_VECTOR_CORE_NUM; int64_t value; aclrtGetDeviceInfo(deviceId, attr, &value); }

此时，可以随时通过aclrtSetDevice接口按线程粒度切换Device（不会影响其他线程）。指定Device后，后续的内存分配、Kernel执行等操作均在该Device上进行，且Stream、Event等也与当前指定的Device相关联。用户可以通过aclrtResetDevice接口释放资源，但更推荐使用aclrtResetDeviceForce接口一次性清理Device上的资源，包括默认Context、默认Stream以及在默认Context下创建的所有Stream。如果默认Context或默认Stream下的任务尚未完成，系统会等待任务完成后才释放资源。在用户程序中，若使用aclrtResetDevice接口，则需确保aclrtSetDevice和aclrtResetDevice接口的调用次数成对出现。

多Device选择的接口调用流程如下图所示：

以下是多Device选择的代码示例，不可以直接拷贝编译运行，仅供参考：

// 指定Device 0作为计算设备，并将Device 0的默认Context作为当前线程的默认Context aclrtSetDevice(0); aclrtStream s0; aclrtCreateStream(&s0); // 执行任务1 ...... // 指定Device 1作为计算设备，并将Device1的默认Context作为当前线程的默认Context aclrtSetDevice(1); aclrtStream s1; aclrtCreateStream(&s1); // 执行任务2 ...... // 复位Device 1，释放计算资源，线程默认Context被释放 // 如需进行继续运行任务，需要显示指定device&context aclrtResetDeviceForce(1); // 复位device0，释放计算资源 aclrtResetDeviceForce(0);

Stream和Event的行为

在与当前Device无所属关系的Stream上下发算子将会失败，示例代码如下：

aclrtSetDevice(0); // 指定Device 0作为计算设备 aclrtStream s0; aclrtCreateStream(&s0); // 在Device 0上创建Stream s0 myKernel<<<8, nullptr, s0>>>(); // 在Device 0上通过Stream s0下发算子 aclrtSetDevice(1); // 指定Device 1作为计算设备 aclrtStream s1; aclrtCreateStream(&s1); // 在Device 1上创建Stream s1 myKernel<<<8, nullptr, s1>>>(); // 在Device 1上通过Stream s1下发算子 // 算子下发失败 myKernel<<<8, nullptr, s0>>>(); // 在Device 1上通过Stream s0下发算子

• 当Stream所属的Device和当前操作的Device不相同时，在此Stream上调用aclrtMemcpyAsync会失败。
• 当Event和Stream关联到不同的Device上时，调用aclrtRecordEvent会失败。
• 当Event和Stream关联到不同的Device上时，调用aclrtStreamWaitEvent会失败。
• 当Event所属的Device和当前操作的Device不相同时，aclrtSynchronizeEvent和aclrtQueryEvent会成功。

跨Device的数据交互

本节中的“跨Device的数据交互”是指一个进程内、根据硬件组网（例如处于PCIe或者HCCS互联的组网拓扑下）、Device之间能够访问彼此的内存。可以使用aclrtDeviceCanAccessPeer接口查询两个Device之间是否支持数据交互，若支持，再根据访问方向，分别调用aclrtDeviceEnablePeerAccess接口开启一个Device到另一个Device的数据交互功能，例如，调用一次aclrtDeviceEnablePeerAccess接口开启Device 0到Device 1的数据交互，再调用一次aclrtDeviceEnablePeerAccess接口开启Device 1到Device 0的数据交互。若需关闭Device之间的数据交互，可调用aclrtDeviceDisablePeerAccess接口。对于两个进程之间的通信请参见进程间通信。

以下是跨Device内存复制的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见Link。

aclInit(NULL); // 初始化 int32_t canAccessPeer = 0; aclrtDeviceCanAccessPeer(&canAccessPeer, 0, 1); // 查询Device 0和Device 1之间是否支持数据交互 if (canAccessPeer == 1) { aclrtSetDevice(0); // Device 0下的操作 uint32_t reserveFlag = 0U; aclrtDeviceEnablePeerAccess(1, reserveFlag); // 开启当前Device(Device 0)到指定Device(Device 1）的数据交互 void *dev0Mem = nullptr; aclrtMalloc(&dev0Mem, 10, ACL_MEM_MALLOC_HUGE_FIRST_P2P); aclrtSetDevice(1); // Device 1下的操作 aclrtDeviceEnablePeerAccess(0, reserveFlag); // 开启当前Device(Device 1)到指定Device(Device 0）的数据交互 void *dev1Mem = nullptr; aclrtMalloc(&dev1Mem, 10, ACL_MEM_MALLOC_HUGE_FIRST_P2P); aclrtMemcpy(dev1Mem, 10, dev0Mem, 10, ACL_MEMCPY_DEVICE_TO_DEVICE); // 将Device 0上的内存数据复制到Device 1上 aclrtDeviceDisablePeerAccess(0); // 关闭当前Device（Device 1）到指定Device（Device 0）的数据交互 aclrtFree(dev1Mem); aclrtResetDevice(1); // 释放Device 1的资源 aclrtSetDevice(0); // 切换到Device 0上进行操作 aclrtDeviceDisablePeerAccess(1); // 关闭当前Device（Device 0）到指定Device（Device 1）的数据交互 aclrtFree(dev0Mem); aclrtResetDeviceForce(0); // Device 0下的操作，调用aclrtResetDevice接口释放Device 0的资源 printf("P2P copy success\n"); } else { printf("current device doesn't support p2p feature\n"); } aclFinalize(); // 去初始化

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime