PyCUDA分阶段共享内存矩阵乘法：无法解释的异常问题

问题背景 #

我实现了一个PyCUDA矩阵乘法核函数（计算C=A*B，其中A为M×K、B为K×N、C为M×N矩阵），包含两项优化：

• 合并全局内存访问
• 分阶段加载共享内存并计算结果
  测试n×n矩阵（n∈[16,32,...,2048]）时出现以下异常：

1. 未添加结果是否超过MAX_VALUE的截断检查时，64×64矩阵结果与CPU计算值随机不符，128×128及更大矩阵则始终错误；
2. 将线程块大小从(32,32)改为(4,4)后，所有测试用例均通过。

疑问1：添加MAX_VALUE截断后64×64用例为何能通过？ #

CPU和GPU的浮点溢出处理行为确实存在差异，你的推测方向是对的：

• x86架构的CPU浮点运算单元（FPU）默认会把溢出的浮点结果截断到对应类型的MAX_VALUE，不会生成无穷大（inf）或NaN这类异常值；而GPU默认的浮点溢出行为是生成inf/NaN，这些异常值会在后续的矩阵乘法累加中扩散，导致结果完全偏离预期。
• 64×64矩阵的计算量较小，只有部分元素会出现溢出，所以结果呈现“随机不符”；128×128及更大矩阵的计算量翻倍增长，溢出元素占比极高，最终导致所有结果都错误。添加截断检查后，GPU端强制将溢出值限制为MAX_VALUE，和CPU的默认行为对齐，因此64×64用例的结果能和CPU保持一致。

疑问2：线程块大小为何会影响计算结果？ #

这和SM资源调度无关——资源不足只会让线程块延迟执行，不会改变计算结果。核心原因是你的核函数中存在线程块内的共享内存访问错误或同步逻辑漏洞：

• 当线程块大小为(32,32)时，每个块包含1024个线程（多数NVIDIA GPU的最大线程块尺寸），如果你的共享内存数组维度、分阶段加载的索引逻辑没有和线程块大小正确匹配，或者__syncthreads()同步点的位置不合理，大线程块会放大这些错误，导致共享内存读写越界、数据未同步就被使用，最终计算出错。
• 换成(4,4)的小线程块后，每个块仅16个线程，错误的索引逻辑刚好在小尺寸下“碰巧”不会触发越界，或者线程间的依赖关系在小线程块下不会暴露同步漏洞，因此计算结果暂时正常。但这只是掩盖了问题，不是真正解决，换其他尺寸的线程块可能还会出问题。

内容的提问来源于stack exchange，提问作者0-0