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 回答