0 overall
Hopper （H100 起）引入的新特性——NVLink SHARP 多播内存（Multicast Memory）机制。
- 通过
cuMulticastCreate创建组播对象，可以让部分内存操作在多个device上被hardware直接broadcast。 - 通过使用
cuMulticastAddDevice讲device加入到组播对象中。 - 每个参与组播的device可以用
cuMulticastBindMem或cuMulticastBindAddr将自己的内存绑定要这个组播对象上。 - 组播的对象可以通过VMM的api（cuMemMap && cuMemSetAccess）映射成某个设备的VA。
1 functions
1.1 cuMulticastGetGranularity
CUresult cuMulticastGetGranularity (size_t* granularity, const CUmulticastObjectProp* prop, CUmulticastGranularity_flags option)
计算指定多播对象属性下的最小或推荐粒度。
- granularity：返回的粒度值。
- prop：多播对象属性的结构体，这里需要填的就是设备数量，multicast需要的memory大小，可导出的handle的bitmask的类型（一般就是FD或者Fabric）。
- option：选择返回最小（MINIMUM）或推荐（RECOMMENDED）粒度。
返回的granularity用来设置multicast对象的大小、绑定offset和地址映射的对齐。拿到granularity大小和prop之后，去下一步的create。
1.2 cuMulticastCreate
CUresult cuMulticastCreate ( CUmemGenericAllocationHandle* mcHandle, const CUmulticastObjectProp* prop)
根据刚刚的prop属性创建一个multicast对象，并返回handle。
- mcHandle：输出参数，返回新创建的多播对象句柄。
- prop：描述多播对象的属性结构体。
- 参与的设备数量由 prop→numDevices 指定。创建后必须通过
cuMulticastAddDevice添加所有设备。 - 绑定内存方法是：cuMulticastBindMem 或 cuMulticastBindAddr。解绑内存使用 cuMulticastUnbind。
- 每个设备可绑定的总内存量由 prop→size 指定。必须是 cuMulticastGetGranularity(…, CU_MULTICAST_GRANULARITY_MINIMUM) 的倍数。或者是对齐到cuMulticastGetGranularity(…, CU_MULTICAST_GRANULARITY_RECOMMENDED) 获得最佳性能。
- 添加完所有设备之后就可以使用VMM的API(cuMemMap,cuMemSetAccess)将multicast映射到VA空间去使用。
- 如果跨进程共享的话就需要去用 cuMemExportToShareableHandle 生成可共享句柄，同时必须在 prop→handleTypes 里指定句柄类型。
- 释放多播对象则使用 cuMemRelease。
1.3 cuMulticastAddDevice
CUresult cuMulticastAddDevice ( CUmemGenericAllocationHandle mcHandle, CUdevice dev )
把某个设备加到已经create完的multicast object内。
- mcHandle：代表多播对象的句柄。
- dev：要加入多播组的设备。 设备总数由 CUmulticastObjectProp::numDevices 在 cuMulticastCreate 时指定。需要注意的是：
- 调用
cuMulticastBindMem或cuMulticastBindAddr之前，必须确保所有设备都已添加； - 同样，调用 cuMemMap（映射多播内存）之前，也必须完成所有设备的添加；
- 若未完成添加，这些函数将会阻塞（等待）。
1.4 cuMulticastBindAddr 和 cuMulticastBindMem
CUresult cuMulticastBindAddr ( CUmemGenericAllocationHandle mcHandle, size_t mcOffset, CUdeviceptr memptr, size_t size, unsigned long long flags )
把一个VA的内存绑定到multicast object上。就是由 memptr 指定的内存区域（必须cudaMallocAsync()或者cuMemCreate分配的内存）绑定到 mcHandle 所表示的多播对象上。
- mcHandle：多播对象句柄。
- mcOffset：多播地址空间中的偏移量。
- memptr：要绑定的内存的虚拟地址。
- size：要绑定的内存大小。
- flags：保留参数，目前必须为 0。
注意点是：
- size、mcOffset、memptr 都必须是由 cuMulticastGetGranularity(…, CU_MULTICAST_GRANULARITY_MINIMUM) 返回的最小粒度的倍数。最佳性能就是把这三个对齐到RECOMMEND返回推荐粒度。
- 绑定的大小不能超过该内存分配本身的大小；
- size + mcOffset 不能超过整个多播对象的大小；
- 绑定的内存必须来自于已加入该多播组的某个设备。
- 如果资源不足，会返回 CUDA_ERROR_OUT_OF_MEMORY；若系统驱动未准备好，则返回 CUDA_ERROR_SYSTEM_NOT_READY；若系统状态非法（驱动守护进程未运行等），会返回 CUDA_ERROR_ILLEGAL_STATE。
CUresult cuMulticastBindMem ( CUmemGenericAllocationHandle mcHandle, size_t mcOffset, CUmemGenericAllocationHandle memHandle, size_t memOffset, size_t size, unsigned long long flags )
与上面的区别就是这里用的不是VA，而是它的handle和在内存里的偏移。
- mcHandle：多播对象句柄。
- mcOffset：在多播对象中的偏移。
- memHandle：表示要绑定的内存的句柄。
- memOffset：在内存中的偏移。
- size：要绑定的内存大小。
- flags：目前必须为 0。
所有偏移量和大小都必须是 cuMulticastGetGranularity 返回的最小粒度的整数倍。 同样，为获得最佳性能，应对齐到推荐粒度。其他的要求与上面一致。
1.5 cuMulticastUnbind
CUresult cuMulticastUnbind ( CUmemGenericAllocationHandle mcHandle, CUdevice dev, size_t mcOffset, size_t size )
unbind指定device上绑定到multicast object的内存区域。
- mcHandle：多播对象句柄。
- dev：包含绑定内存的设备。
- mcOffset：在多播对象中的偏移量。
- size：要解绑的大小。
解除绑定的偏移与大小必须是 cuMulticastGetGranularity(…, MINIMUM) 的倍数。 size + mcOffset 不能超过多播对象总大小。mcOffset 和 size 必须与绑定时使用的值完全一致，否则行为未定义。