1

我有以下问题。我正在尝试将共享数组划分为更小的数组，然后在其他设备函数中使用这些数组。在我的内核函数中，

for (int block_x = 0; block_x < blockDim.x; block_x++) {
  for (int block_y = 0; block_y < blockDim.y; block_y++) {
  //set up shared memory block
  extern __shared__ vec3f share[];
  vec3f *sh_pos = share;
  vec3f *sh_velocity = &sh_pos[blockDim.x*blockDim.y];
  vec3f *sh_density = &sh_velocity[blockDim.x*blockDim.y];
  vec3f *sh_pressure = &sh_density[blockDim.x*blockDim.y];
  //index by 2d threadidx's
  unsigned int index = (block_x * blockDim.x + threadIdx.x) + blockDim.x * gridDim.x * (block_y * blockDim.y + threadIdx.y);
  sh_pos[blockDim.x * threadIdx.x + threadIdx.y] = oldParticles[index].position();
  sh_velocity[blockDim.x * threadIdx.x + threadIdx.y] = oldParticles[index].velocity();
  sh_pressure[blockDim.x * threadIdx.x + threadIdx.y].x = oldParticles[index].pressure();
  sh_density[blockDim.x * threadIdx.x + threadIdx.y].x = oldParticles[index].density();
  __syncthreads();
  d_force_pressure(oldParticles[arr_pos],c_kernel_support);
  __syncthreads();
  }
}

据我所知，所有“sh_”数组都填充了零，而不是我想要的值。我不能说我做错了什么。请注意， vec3f 是浮点向量，就像 float3 数据类型一样。另外，我不认为我可以混合使用密度和压力的浮点数，所以我只是将它们制作成矢量并使用单个组件。然后，例如我的 d_force_pressure 函数是，

__device__ void d_force_pressure(particle& d_particle, float h) {
  extern __shared__ vec3f share[];
  vec3f *sh_pos = share;
  vec3f *sh_velocity = &sh_pos[blockDim.x*blockDim.y];
  vec3f *sh_density = &sh_velocity[blockDim.x*blockDim.y];
  vec3f *sh_pressure = &sh_density[blockDim.x*blockDim.y];
  for (int i = 0; i < blockDim.x * blockDim.y; i++) {
    vec3f diffPos = d_particle.position() - sh_pos[i];
    d_particle.force() += GradFuncion(diffPos,h) * -1.0 * c_particle_mass *  (d_particle.pressure()+sh_pressure[i].x)/(2.0*sh_density[i].x);
  }  
 }

调用此函数后，我得到 NaN，因为我除以零（sh_density[i].x据我所知，为 0）。这通常也是加载共享内存的正确方法吗？

内核被调用

dim3 block(BLOCK_SIZE,BLOCK_SIZE,1);
dim3 grid((int)ceil(sqrt(float(max_particles)) / (float(block.x*block.y))), (int)ceil(sqrt(float(max_particles)) / (float(block.x*block.y))), 1);
int sharedMemSize = block.x*block.y*4*sizeof(vec3f);
force_kernel<<< grid,block,sharedMemSize  >>>(particle_ptrs[1],particle_ptrs[0],time_step);
4

2 回答 2

1

这是一种后续的答案。

根据@RobertCrovella 的评论，我继续运行 cuda-memcheck。信不信由你，这实际上没有显示任何错误。但是，当我更改代码中的常量（控制某些数组的大小）时，cuda-memcheck 显示与此处发布的问题相关的错误 write error。这让我重新检查了填充共享数组的方式。基本上需要改变的是

for (int block_x = 0; block_x < blockDim.x; block_x++) {
  for (int block_y = 0; block_y < blockDim.y; block_y++) {

至

for (int block_x = 0; block_x < gridDim.x; block_x++) {
  for (int block_y = 0; block_y < gridDim.y; block_y++) {

我相信这会为index变量提供正确的位置。我基本上了解到，每当您使用共享内存并注意到运行缓慢时，使用 cuda-memcheck 是个好主意。

于 2013-05-26T05:25:02.223 回答
0

我在您之前的问题中表示您不想这样做：

dim3 grid((int)ceil(sqrt(float(max_particles)) / (float(block.x*block.y))), (int)ceil(sqrt(float(max_particles)) / (float(block.x*block.y))), 1);

你想这样做：

dim3 grid((int)ceil(sqrt(float(max_particles)) / (float(block.x))), (int)ceil(sqrt(float(max_particles)) / (float(block.y))), 1);

x 网格方向应按线程块 x 维度缩放，而不是线程块 x 维度 * 线程块 y 维度。但是我在之前的答案中发布的代码也有这个错误，即使我在评论中指出了它，我也忘了修复它。

此外，这个索引对我来说看起来不正确：

sh_velocity[blockDim.x * threadIdx.x + threadIdx.y]

我认为应该是：

sh_velocity[blockDim.x * threadIdx.y + threadIdx.x]

你有几个例子。

您尚未发布完整的可执行文件。当然，可能存在比我上面指出的问题更多的问题。如果我必须完成我在上一个问题中所做的所有 vec3f -> float3 转换工作，那么其他人可以帮助你。如果您编写一个不依赖于我没有的一堆代码的简单复制器，我可以尝试进一步提供帮助。如果您这样做，您很可能会自己发现问题。

您是否像我在上一个答案中建议的那样将 cuda 错误检查放入您的代码中？

您可能还想通过 cuda-memcheck 运行您的代码：

cuda-memcheck ./mycode
于 2013-05-25T14:30:00.057 回答