CUDA粒子中的最近邻

编辑2: 请看一下 this crosspost TLDR。

编辑 :假设粒子被分割成网格单元(例如 16^3

在这种情况下,我可以将相邻单元中的所有粒子加载到本地内存中,并通过迭代来计算某些属性。然后我可以将特定值写入当前网格单元中的每个粒子。

这种方法是否有利于对所有粒子和每个迭代(大多数情况下是相同的)邻居运行内核?

number of particles/number of grid cells ?

我正在尝试重新实现(和修改) CUDA Particles 并使用它来查询每个粒子的最近邻居。我创建了以下结构:

• 缓冲器 P 保持所有粒子的3D位置( float3 )

• 缓冲器 Sp 存储 int2 粒子id对及其空间散列。 服务提供商 根据哈希排序。(散列只是从3D到1D的简单线性映射,还没有Z索引。)

• 缓冲器 L 存储 整数2 缓冲区中特定空间哈希的起始和结束位置对 服务提供商 。示例: L[12] = (int2)(0, 50) .

  • L[12].x 是索引(in 服务提供商 第一 带空间散列的粒子 12 .
  • L[12].y 是索引(in 服务提供商 )的 最后的 带空间散列的粒子 12 .

现在我有了所有这些缓冲区,我想迭代所有粒子 P 对于每个粒子,迭代其最近的邻居。目前我有一个内核,看起来像这样(伪代码):

__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
  size_t gid             = get_global_id(0);
  float3 curr_particle   = P[gid];
  int    processed_value = 0;

  for(int x=-1; x<=1; x++)
    for(int y=-1; y<=1; y++)
      for(int z=-1; z<=1; z++) {

        float3 neigh_position = curr_particle + (float3)(x,y,z)*GRID_CELL_SIDE;

        // ugly boundary checking
        if ( dot(neigh_position<0,        (float3)(1)) +
             dot(neigh_position>BOUNDARY, (float3)(1))   != 0)
             continue;

        int neigh_hash        = spatial_hash( neigh_position );
        int2 particles_range  = L[ neigh_hash ];

        for(int p=particles_range.x; p<particles_range.y; p++)
          processed_value += heavy_computation( P[ Sp[p].y ] );

      }

  Out[gid] = processed_value;
}

这段代码的问题是速度慢。我怀疑非线性GPU内存访问(特别是 P[Sp[p].y] 在内心深处 for

我想做的是使用 Z-order curve 作为空间散列。那样我就只能吃一个了 对于 当查询邻居时,循环遍历连续的内存范围。唯一的问题是,我不知道应该是什么开始和停止Z索引值。

我想要实现的圣杯:

__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
  size_t gid             = get_global_id(0);
  float3 curr_particle   = P[gid];
  int    processed_value = 0;

  // How to accomplish this??
  // `get_neighbors_range()` returns start and end Z-index values
  // representing the start and end near neighbors cells range
  int2 nearest_neighboring_cells_range = get_neighbors_range(curr_particle);
  int first_particle_id = L[ nearest_neighboring_cells_range.x ].x;
  int last_particle_id  = L[ nearest_neighboring_cells_range.y ].y;

  for(int p=first_particle_id; p<=last_particle_id; p++) {
      processed_value += heavy_computation( P[ Sp[p].y ] );
  }

  Out[gid] = processed_value;
}