位与操作是否比<int>强制类型转换更高效？CUDA代码优化疑问

原代码片段 #

研究的是NVIDIA CUDA示例2_Concepts_and_Techniques/particles中的以下代码：

// calculate position in uniform grid
__device__ int3 calcGridPos(float3 p) {
  int3 gridPos;
  gridPos.x = floorf((p.x - params.worldOrigin.x) / params.cellSize.x);
  gridPos.y = floorf((p.y - params.worldOrigin.y) / params.cellSize.y);
  gridPos.z = floorf((p.z - params.worldOrigin.z) / params.cellSize.z);
  return gridPos;
}

// calculate address in grid from position (clamping to edges)
__device__ uint calcGridHash(int3 gridPos) {
  gridPos.x = gridPos.x &
              (params.gridSize.x - 1);  // wrap grid, assumes size is power of 2
  gridPos.y = gridPos.y & (params.gridSize.y - 1);
  gridPos.z = gridPos.z & (params.gridSize.z - 1);
  return __umul24(__umul24(gridPos.z, params.gridSize.y), params.gridSize.x) +
         __umul24(gridPos.y, params.gridSize.x) + gridPos.x;
}

用户疑问 #

1. calcGridPos函数将floorf（返回float类型）的结果存入int3类型变量，推测这和后续函数相关，该理解是否正确？
2. calcGridHash中的位与操作gridPos.x & (params.gridSize.x - 1)，猜测是为了高效完成float转int，同时配合__umul24优化乘法，该理解是否正确？若替换为以下代码是否安全？

// calculate address in grid from position (clamping to edges)
__device__ uint calcGridHash(int3 gridPos) {
  gridPos.x = (int)gridPos.x;
  gridPos.y = (int)gridPos.y;
  gridPos.z = (int)gridPos.z;
  return ((gridPos.z * params.gridSize.y) * params.gridSize.x) +
         (gridPos.y * params.gridSize.x) + gridPos.x;
}

解答 #

关于calcGridPos的类型赋值 #

你的推测方向正确，这段代码完全合法：

• floorf返回的float值本身是整数形式（比如floorf(3.7)得到3.0，floorf(-2.1)得到-3.0），将其直接赋值给int类型变量时，CUDA编译器会自动完成浮点到整数的截断转换，结果和取整一致。
• 后续calcGridHash需要整数类型的网格坐标计算哈希值，所以这里的类型转换是服务于后续计算的，官方写法没有问题。

关于位与操作的作用及替换代码的安全性 #

你的理解有误，位与操作的核心作用不是转int，而是网格坐标的环绕处理：

• 注释明确说明网格大小是2的幂，此时params.gridSize.x -1的二进制是全1形式（比如大小为64时，减1是63，二进制为0b111111），位与操作等价于gridPos.x % params.gridSize.x，但位运算执行速度远快于模运算。
• 这个操作的目的是当粒子超出网格边界时，将坐标“包裹”到网格内部，避免哈希值超出有效范围。

替换后的代码不安全，原因如下：

1. gridPos.x = (int)gridPos.x完全多余，因为gridPos本身就是int3类型，无需再次转换。
2. 去掉位与操作后，若粒子的网格坐标超出[0, params.gridSize.x-1]范围，计算出的哈希值会超出网格哈希表的有效索引，直接引发数组越界访问，导致内存错误或计算结果混乱。
3. 替换__umul24为普通乘法会损失性能（__umul24是CUDA针对24位整数优化的乘法指令，在多数架构上更快），但只要计算结果不超出uint的范围，不会直接引发错误。

内容的提问来源于stack exchange，提问作者Sangjun Lee