本文对应 MXMACA 编程接口第 3 章中的3.1 设备管理与3.1.1 查询设备信息。目标不是死记 API而是理解程序如何发现 GPU、读取 GPU 的硬件能力并据此决定 Kernel核函数如何配置。一、为什么要先学习“设备管理”一台服务器中可能安装多张 GPU例如GPU 0 GPU 1 GPU 2 GPU 3程序在启动 Kernel核函数之前至少需要知道当前系统有多少张可用 GPU每张 GPU 的名称和硬件能力是什么每个 Block线程块最多可以放多少线程一个 Wave波前包含多少线程每个 Block 可以使用多少 Shared Memory共享内存和 Register寄存器Grid网格的最大维度是多少GPU 是否支持数据传输和计算重叠这些信息会直接影响Grid 怎么划分 Block 配置多大 Tile 应该切多大 每个 Block 能使用多少共享内存 寄存器会不会使用过多 Kernel 能不能充分占满 GPU因此设备管理不是单纯地“打印显卡名称”而是后续性能优化的硬件基础。二、设备管理的整体思路MXMACA 提供了一组 Runtime API运行时接口供 Host主机端程序查询和管理 GPU。本节最重要的两个函数是mcGetDeviceCount(count);mcGetDeviceProperties(prop,device_id);它们分别负责API中文含义作用mcGetDeviceCount获取设备数量查询当前系统中有多少张可用 GPUmcGetDeviceProperties获取设备属性查询指定 GPU 的硬件参数mcDeviceProp_t设备属性结构体保存查询到的设备信息程序执行流程Host 程序启动 ↓ 查询 GPU 数量 ↓ 遍历每一张 GPU ↓ 查询该 GPU 的属性 ↓ 打印或保存硬件信息 ↓ 根据硬件能力配置 Kernel三、查询 GPU 数量mcGetDeviceCount基本写法intcount;mcGetDeviceCount(count);1.count是什么count用来保存系统检测到的 GPU 数量。例如系统中有 4 张 GPUcount 4设备编号通常是device 0 device 1 device 2 device 32. 为什么参数是countcount表示把变量count的地址传给函数。可以理解为程序这里有一个空盒子 count 程序把盒子的地址交给 mcGetDeviceCount 函数检测 GPU 数量 函数把结果写进这个盒子调用结束后程序就可以读取count的值。四、查询设备属性mcGetDeviceProperties基本写法mcDeviceProp_t prop;mcGetDeviceProperties(prop,i);其中参数含义prop接收设备属性的结构体地址i要查询的 GPU 编号假设i0;那么程序查询的是device 0结果写入prop。后续可以访问prop.name prop.waveSize prop.totalGlobalMem prop.maxThreadsPerBlock五、完整示例代码#includestdio.h#includemaca_runtime.hintmain(void){mcDeviceProp_t prop;intcount;mcGetDeviceCount(count);for(inti0;icount;i){mcGetDeviceProperties(prop,i);printf(--- General Information for device %d ---\n,i);printf(Name: %s\n,prop.name);printf(Compute capability: %d.%d\n,prop.major,prop.minor);printf(Clock rate: %ld\n,prop.clockRate);printf(Device copy overlap: );if(prop.deviceOverlap)printf(Enabled\n);elseprintf(Disabled\n);printf(Kernel execution timeout: );if(prop.kernelExecTimeoutEnabled)printf(Enabled\n);elseprintf(Disabled\n);printf(--- Memory Information for device %d ---\n,i);printf(Total global mem: %ld\n,prop.totalGlobalMem);printf(Total constant mem: %ld\n,prop.totalConstMem);printf(Max mem pitch: %ld\n,prop.memPitch);printf(Texture Alignment: %ld\n,prop.textureAlignment);printf(--- Multi-Processor Information for device %d ---\n,i);printf(Multi-Processor count: %d\n,prop.multiProcessorCount);printf(Shared mem per block: %ld\n,prop.sharedMemPerBlock);printf(Registers per block: %d\n,prop.regsPerBlock);printf(Threads in wave: %d\n,prop.waveSize);printf(Max threads per block: %d\n,prop.maxThreadsPerBlock);printf(Max thread dimensions: (%d, %d, %d)\n,prop.maxThreadsDim[0],prop.maxThreadsDim[1],prop.maxThreadsDim[2]);printf(Max grid dimensions: (%d, %d, %d)\n,prop.maxGridSize[0],prop.maxGridSize[1],prop.maxGridSize[2]);printf(\n);}return0;}头文件名称应以实际 MXMACA 开发环境为准。手册示例可能已经在工程配置中统一包含了对应头文件。六、逐段理解示例代码6.1 定义设备属性结构体mcDeviceProp_t prop;mcDeviceProp_t可以理解为一张“GPU 参数表”其中保存设备名称 计算能力 时钟频率 显存容量 共享内存容量 寄存器数量 Wave 大小 Block 线程上限 Grid 维度上限6.2 获取设备数量intcount;mcGetDeviceCount(count);执行后count 当前可用 GPU 数量6.3 遍历每一张 GPUfor(inti0;icount;i){...}如果count 4循环依次查询设备 0、1、2、3。6.4 获取第i张 GPU 的属性mcGetDeviceProperties(prop,i);执行后第i张设备的信息被写入prop。七、输出字段具体代表什么7.1Name设备名称prop.name示例Name: Device 4000设备名称主要用于确认程序运行在哪种 GPU 上在多型号服务器中区分设备保存 Benchmark性能测试环境信息按架构选择不同 Kernel 实现。7.2Compute capability计算能力版本prop.major prop.minor示例Compute capability: 10.2它通常表示设备架构或运行时支持能力的版本。需要注意这里虽然使用了类似 CUDA 的Compute capability计算能力命名但不同厂商的版本体系不一定可以直接横向比较。7.3Clock rate时钟频率prop.clockRate示例Clock rate: 1600000具体单位应以 MXMACA 手册定义为准常见运行时接口可能使用 kHz千赫兹。时钟频率高不代表 Kernel 一定快性能还受以下因素影响内存带宽 计算单元数量 并行度 寄存器使用量 共享内存使用量 访存是否连续 是否存在分支发散7.4Device copy overlap复制与计算重叠能力prop.deviceOverlap示例Device copy overlap: Enabled表示设备是否具备数据传输和 Kernel 计算重叠执行的能力。时间轴 → 数据传输████████████████ Kernel ████████████████但Enabled只代表硬件具备能力不代表程序自动实现重叠。通常还要配合Asynchronous Copy异步复制Stream流合理的任务调度满足要求的 Host Memory主机内存配置。7.5Kernel execution timeoutKernel 执行超时prop.kernelExecTimeoutEnabled如果该功能开启执行时间过长的 Kernel 可能被系统终止。服务器计算卡通常更希望Kernel execution timeout Disabled因为深度学习和高性能计算 Kernel 可能执行较长时间。八、内存相关属性8.1Total global mem全局显存总量prop.totalGlobalMem示例Total global mem: 68283269120若单位为 Byte字节按 GiB 换算68283269120 ÷ 1024³ ≈ 63.59 GiBGlobal Memory全局内存主要保存输入 Tensor 权重 Tensor 输出 Tensor 中间结果在 Fused MoE GEMM 中需要考虑X_e 是否放得下 W_e 是否放得下 多个 Expert 权重总量 临时 Buffer 占用显存容量足够不代表性能高性能还要看带宽、访问次数和数据复用。8.2Total constant mem常量内存总量prop.totalConstMemConstant Memory常量内存通常用于只读数据容量较小适合多个线程读取相同内容。8.3Max mem pitchPitch 内存最大跨度prop.memPitchPitch行跨度用于二维内存分配表示相邻两行起始地址之间的距离。实际数据长度12 Byte 底层行跨度 16 Byte其中 16 Byte 就可以理解为 Pitch。8.4Texture Alignment纹理内存对齐prop.textureAlignment表示纹理或相关内存资源的地址对齐要求。对当前 TileLang 和 GEMM 入门不是第一优先级先理解概念即可。九、计算单元与并行相关属性9.1Multi-Processor count多处理器数量prop.multiProcessorCount示例Multi-Processor count: 104可以把 Multi-Processor多处理器MP理解为 GPU 上执行线程块的重要计算单元。GPU ├── MP 0 ├── MP 1 ├── ... └── MP 103如果 Grid 中只有 4 个 Block而设备有 104 个 MP大量 MP 可能没有工作。对于 Fused MoE GEMM以下情况可能导致 Block 数量不足Expert 数量较少每个 Expert 的M_e很小Tile 设置过大。9.2Shared mem per block每个 Block 的共享内存上限prop.sharedMemPerBlock示例Shared mem per block: 65536若单位为 Byte65536 Byte 64 KiBGEMM 常见数据流Global Memory ↓ Shared Memory ↓ Register ↓ 矩阵乘加GEMM 中 Tile 的共享内存占用该怎么计算在 GEMMGeneral Matrix Multiplication通用矩阵乘法优化中经常会看到下面这些参数BLOCK_M 64 BLOCK_N 64 BLOCK_K 32 float32 4 Byte这些参数不是线程数量而是描述矩阵分块也就是 Tile数据块的大小。1. BLOCK_M、BLOCK_N、BLOCK_K 分别是什么意思假设矩阵乘法为C A × B矩阵形状为A: [M, K] B: [K, N] C: [M, N]参数含义如下参数中文含义作用BLOCK_MM 方向分块大小一次计算输出矩阵 C 的多少行BLOCK_NN 方向分块大小一次计算输出矩阵 C 的多少列BLOCK_KK 方向分块大小每次沿 K 方向处理多少个元素假设BLOCK_M 64 BLOCK_N 64 BLOCK_K 32那么一个 Block线程块通常负责计算一个C_tile: [64, 64]也就是输出矩阵 C 中的一块64 × 64区域。2. 为什么需要 A_tile 和 B_tile为了计算一个64 × 64的C_tile不能直接一次完成全部 K 方向的计算。通常会沿 K 方向分段处理。当前设置BLOCK_K 32表示每次只处理 K 方向的 32 个元素。因此每一轮需要加载A_tile: [BLOCK_M, BLOCK_K] B_tile: [BLOCK_K, BLOCK_N]代入参数A_tile: [64, 32] B_tile: [32, 64]矩阵形状对应关系如下A_tile B_tile K 32 N 64 ┌──────────────┐ ┌──────────────────────┐ │ │ │ │ │ │ │ │ │ │ M64 │ │ K32 │ │ │ │ └──────────────┘ └──────────────────────┘ A_tile × B_tile ↓ C_tile ┌────────────────┐ │ │ │ 64 × 64 │ │ │ └────────────────┘也就是[64 × 32] × [32 × 64] ↓ [64 × 64]3. A_tile 为什么占用 8192 ByteA_tile的形状为64 × 32元素数量为64 × 32 2048 个元素因为数据类型是float321 个 float32 4 Byte所以A_tile 占用空间 64 × 32 × 4 8192 Byte换算成 KiB8192 ÷ 1024 8 KiB因此A_tile 8192 Byte 8 KiB4. B_tile 为什么也占用 8192 ByteB_tile的形状为32 × 64元素数量为32 × 64 2048 个元素因为每个float32占 4 ByteB_tile 占用空间 32 × 64 × 4 8192 Byte 8 KiB因此B_tile 8192 Byte 8 KiB5. 为什么共享内存总需求是 16 KiB一个 Block 在计算当前C_tile时需要同时缓存A_tile 8 KiB B_tile 8 KiB所以共享内存总需求为Shared Memory A_tile B_tile 8 KiB 8 KiB 16 KiB6. 为什么这里没有计算 C_tile 的共享内存这里计算的 16 KiB 只包含A_tile B_tile没有包含C_tile。原因是高性能 GEMM 中输出矩阵 C 的中间结果通常保存在 Register寄存器里而不是完整地存放在 Shared Memory共享内存中。典型的数据流如下A、BGlobal Memory全局内存 ↓ A_tile、B_tileShared Memory共享内存 ↓ 线程读取小块数据 ↓ C 的部分结果Register寄存器 ↓ 最终写回 Global Memory全局内存因此共享内存主要缓存 A_tile 和 B_tile 寄存器主要保存 C 的累加结果不过一些复杂 Kernel 在数据重排、转置或者 Epilogue尾处理阶段也可能额外使用共享内存。所以这里计算的是一个近似的基础需求。7. 完整 K 很大时为什么不乘以循环次数假设完整矩阵的 K 为K 256而BLOCK_K 32那么需要处理256 ÷ 32 8 轮计算过程如下第 1 轮 加载 A[:, 0:32] 加载 B[0:32, :] 执行乘法并累加 第 2 轮 加载 A[:, 32:64] 加载 B[32:64, :] 执行乘法并累加 第 3 轮 加载 A[:, 64:96] 加载 B[64:96, :] 执行乘法并累加 ... 第 8 轮 加载最后一段数据 执行乘法并累加共享内存中的空间会被反复复用。也就是说第 1 轮使用这 16 KiB 第 2 轮覆盖并继续使用这 16 KiB 第 3 轮继续复用这 16 KiB因此单缓冲情况下共享内存仍然只需要16 KiB而不是16 KiB × 8这体现了共享内存的核心作用把 Global Memory 中的数据按 Tile 加载到片上存储中并循环复用这块高速内存。8. 什么是双缓冲为了让数据加载和矩阵计算重叠执行Kernel 可能使用 Double Buffering双缓冲。单缓冲过程加载数据 ↓ 等待加载完成 ↓ 执行计算 ↓ 继续加载下一轮数据双缓冲会准备两套共享内存空间Buffer 0当前正在计算的数据 Buffer 1提前加载下一轮数据执行过程可以理解为时间轴 → Buffer 0加载 → 计算 → 加载 → 计算 Buffer 1 加载 → 计算 → 加载 → 计算这样可以让当前 Tile 正在计算 下一轮 Tile 同时加载从而减少等待 Global Memory 的时间。如果单缓冲需要16 KiB那么双缓冲通常近似需要16 KiB × 2 32 KiB即两份 A_tile 两份 B_tile9. 与设备共享内存上限有什么关系假设通过设备信息查询得到sharedMemPerBlock 65536 Byte换算为65536 Byte 64 KiB当前 Tile 配置的单缓冲需求16 KiB 64 KiB从容量上看可以正常使用。双缓冲需求32 KiB 64 KiB从容量上看也可能可以使用。但需要注意没有超过共享内存上限只代表配置合法不代表性能一定最好。因为一个 MPMulti-Processor多处理器上的共享内存资源是有限的。假设一个 Block 使用32 KiB Shared Memory而一个 MP 可用于运行 Block 的共享内存大约为64 KiB那么同一个 MP 上最多可能同时驻留64 KiB ÷ 32 KiB 2 个 Block如果每个 Block 使用60 KiB那么一个 MP 上可能只能驻留 1 个 Block。这可能导致同时运行的 Block 数量减少 ↓ Occupancy占用率降低 ↓ 隐藏访存延迟的能力下降因此共享内存使用需要在以下两者之间进行平衡更大的 Tile → 更高的数据复用率 但同时 更大的 Tile → 更高的共享内存占用 → 可能降低 Occupancy10. 修改 BLOCK_N 后会发生什么假设修改为BLOCK_M 64 BLOCK_N 128 BLOCK_K 32 float32 4 Byte那么A_tile 64 × 32 × 4 8192 Byte 8 KiBB_tile 32 × 128 × 4 16384 Byte 16 KiB共享内存总需求8 KiB 16 KiB 24 KiB如果使用双缓冲24 KiB × 2 48 KiB虽然仍然低于 64 KiB但共享内存压力已经明显增加。11. 如果数据类型是 float16 呢float16每个元素只占2 Byte仍然使用BLOCK_M 64 BLOCK_N 64 BLOCK_K 32则A_tile 64 × 32 × 2 4096 Byte 4 KiBB_tile 32 × 64 × 2 4096 Byte 4 KiB总共享内存4 KiB 4 KiB 8 KiB和float32相比float3216 KiB float168 KiB这说明数据类型会直接影响Tile 能切多大 每个 Block 的共享内存占用 一个 MP 能驻留多少个 Block Kernel 的 Occupancy12. 与 TileLang 和 Fused MoE GEMM 的关系在 TileLang Kernel 中可能会看到类似配置BLOCK_M64BLOCK_N64BLOCK_K32这些参数会决定一个 Block 负责多大的输出 Tile 每次从 Global Memory 搬多少数据 共享内存需要多大 寄存器需要保存多少累加结果 Grid 中会生成多少个 Block在 Fused MoE GEMM 中每个 Expert专家执行Y_e X_e × W_e其中X_e: [M_e, K] W_e: [K, N] Y_e: [M_e, N]这里的M_e由该专家收到的 Token 数决定假设BLOCK_M 64 BLOCK_N 64那么一个 Block 通常负责64 个 Token × 64 个输出维度如果某个 Expert 只有 20 个 TokenM_e 20但BLOCK_M 64那么一个 Tile 中可能只有 20 行有效数据其余位置需要通过边界判断屏蔽。这会造成一定资源浪费。因此在比赛中需要同时考虑Tile 大小 共享内存占用 寄存器占用 Wave 利用率 Expert 的 Token 分布 Grid 中 Block 数量13. 最短记忆方法可以把整个过程记成一个 Block 要计算一个 64 × 64 的 C_tile。 为了完成一轮 BLOCK_K 32 的计算需要加载 A_tile64 × 32 B_tile32 × 64 float32 每个元素占 4 Byte。 所以共享内存需求为 64 × 32 × 4 32 × 64 × 4 16384 Byte 16 KiB一句话总结一个 Block 在处理一轮BLOCK_K 32的矩阵乘法时需要把A_tile和B_tile缓存到 Shared Memory 中当BLOCK_M 64、BLOCK_N 64、数据类型为float32时共享内存需求约为 16 KiB。若使用 Double Buffering双缓冲16 KiB × 2 32 KiB仍需考虑额外开销。手册截图中的打印文字可能写成 “Shared mem per MP”但代码读取的是sharedMemPerBlock。应以字段定义为准不要把“每个 MP”和“每个 Block”混为一谈。9.3Registers per block每个 Block 的寄存器数量prop.regsPerBlockRegister寄存器是线程私有、速度非常快的片上存储常用于保存累加器 accumulator 线程私有中间结果 索引变量 临时数据寄存器使用过多可能导致每线程寄存器增加 ↓ 一个 MP 可同时驻留的线程减少 ↓ Occupancy占用率下降 ↓ 隐藏访存延迟的能力下降因此需要平衡更多寄存器复用 VS 更高并行驻留能力9.4Threads in wave一个 Wave 的线程数prop.waveSize示例Threads in wave: 64即1 Wave 64 Threads因此 Block 大小通常优先考虑 64 的倍数64 128 256 512例如Block 128 Threads 128 ÷ 64 2 Waves如果Block 96 Threads可能需要两个 Wave第 1 个 Wave64 个有效线程 第 2 个 Wave32 个有效线程 32 个空闲位置这可能造成资源浪费。但 Block 大小是否合适还要结合寄存器、共享内存、访存和计算量判断。9.5Max threads per block每个 Block 的最大线程数prop.maxThreadsPerBlock示例Max threads per block: 1024必须满足blockDim.x × blockDim.y × blockDim.z ≤ maxThreadsPerBlock例如dim3block(32,32);总线程数为 1024达到上限。而dim3block(64,32);总线程数为 2048超过上限Kernel 无法正常启动。9.6Max thread dimensionsBlock 各维度上限prop.maxThreadsDim[0]prop.maxThreadsDim[1]prop.maxThreadsDim[2]分别对应blockDim.x 上限 blockDim.y 上限 blockDim.z 上限容易误解的是每个维度都不超过 1024并不代表三个维度可以同时取 1024。仍必须满足总线程数不超过maxThreadsPerBlock。9.7Max grid dimensionsGrid 各维度上限prop.maxGridSize[0]prop.maxGridSize[1]prop.maxGridSize[2]分别对应gridDim.x gridDim.y gridDim.z二维矩阵计算中常见dim3grid(ceil_div(N,BLOCK_N),ceil_div(M,BLOCK_M));其中grid.x 列方向 Tile 数量 grid.y 行方向 Tile 数量十、设备属性如何影响 TileLang 和比赛优化比赛核心是 Fused MoE GEMMY_e X_e W_e其中X_e: [M_e, K] W_e: [K, N] Y_e: [M_e, N]设备信息和优化参数的对应关系设备属性影响的优化决策waveSizenum_threads是否适合 Wave 大小maxThreadsPerBlock每个 Block 最多使用多少线程sharedMemPerBlockBLOCK_M、BLOCK_N、BLOCK_K能切多大regsPerBlockRegister Blocking寄存器分块规模multiProcessorCountGrid 需要多少 Block 才能占满设备totalGlobalMem输入、权重和临时 Buffer 是否放得下deviceOverlap是否值得设计传输与计算重叠maxGridSize大规模任务的 Grid 是否合法例如 TileLang Kernel 可能配置BLOCK_M64BLOCK_N128BLOCK_K32num_threads128需要检查128 是否与 waveSize64 匹配 共享内存是否超过 sharedMemPerBlock 寄存器使用是否过高 Grid 中 Block 数量是否足以覆盖 104 个 MP 小 Expert 会不会因为 Tile 太大而浪费计算资源所以设备查询的真正作用是把“拍脑袋调参数”变成“基于硬件约束调参数”。十一、结合 Fused MoE GEMM 的具体例子假设BLOCK_M 64 BLOCK_N 128 BLOCK_K 32 num_threads 128 float16 2 Byte共享内存近似需求A_tile 64 × 32 × 2 4096 Byte B_tile 32 × 128 × 2 8192 Byte 合计 12288 Byte 12 KiB启用双缓冲后12 KiB × 2 24 KiB若每个 Block 支持 64 KiB 共享内存则容量方面可行。线程方面num_threads 128 waveSize 64所以一个 Block 包含 2 个 Wave。若某 ExpertM_e 20 BLOCK_M 64那么只处理 20 个 Token却按 64 行 Tile 设计可能有大量无效位置。这说明大 Tile 有利于数据复用 但可能不适合小 Expert设备属性告诉你“能不能这样配”性能测试才告诉你“这样配快不快”。十二、建议增加错误检查正式工程中建议检查 Runtime API 返回值。下面为伪代码intcount0;mcError_t errmcGetDeviceCount(count);if(err!mcSuccess){printf(mcGetDeviceCount failed\n);return-1;}for(inti0;icount;i){mcDeviceProp_t prop;errmcGetDeviceProperties(prop,i);if(err!mcSuccess){printf(Failed to query device %d\n,i);continue;}printf(Device %d: %s\n,i,prop.name);}mcError_t和成功状态常量的具体名称应以当前版本 MXMACA Runtime API 文档为准。还应处理count 0即没有检测到可用 GPU 的情况。十三、常见误区误区 1GPU 数量多单个 Kernel 就一定快错误。GPU 数量影响多卡能力但单个 Kernel 默认通常只在一张 GPU 上执行。误区 2Block 线程数越大越好错误。线程数过大可能导致寄存器压力、共享内存压力和 Occupancy 下降。误区 3共享内存没超过上限就一定性能好错误。还要考虑访存冲突、数据复用率、寄存器、同步开销和计算访存比。误区 4waveSize 64所有 Block 都必须是 64 个线程错误。更准确地说Block 线程数通常优先选择 64 的倍数但要综合评估。误区 5最大线程维度是(1024,1024,1024)Block 就能有1024³个线程错误。还必须满足三个维度乘积不超过maxThreadsPerBlock。误区 6支持 Copy/Compute Overlap程序就会自动重叠错误。通常还需要异步 API、Stream 和合理流水线。十四、本节知识结构总结设备管理 │ ├── mcGetDeviceCount │ └── 查询可用 GPU 数量 │ ├── mcGetDeviceProperties │ └── 查询指定 GPU 属性 │ └── mcDeviceProp_t ├── 设备名称与架构 ├── 时钟频率 ├── 显存容量 ├── 共享内存 ├── 寄存器 ├── MP 数量 ├── Wave 大小 ├── Block 上限 └── Grid 上限十五、一句话总结mcGetDeviceCount告诉程序“有几张 GPU”mcGetDeviceProperties告诉程序“每张 GPU 能做什么”Kernel 优化就是在这些硬件限制内合理设计 Grid、Block、Wave、Tile、Shared Memory 和 Register。十六、与后续学习的衔接下一步通常会继续学习选择当前设备 ↓ 设备内存分配 ↓ Host 与 Device 数据传输 ↓ Kernel 启动 ↓ 同步与错误检查完整 GPU 程序流程查询设备 → 选择设备 → 分配内存 → 传输数据 → 启动 Kernel → 同步 → 获取结果 → 释放资源