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 データ型と同様に float のベクトルであることに注意してください。また、密度と圧力のためにフロートを混ぜることができるとは思わなかったので、それらをベクトルにして単一のコンポーネントを使用しています。次に、たとえば、私の 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 は、ここに投稿された書き込みエラーに関連するエラーを示しました。これにより、共有配列を埋める方法を再確認しました。基本的に変更が必要だったのは

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 寸法 * ねじブロックの y 寸法ではなく、ねじブロックの x 寸法でスケーリングする必要があります。ただし、以前の回答で投稿したコードにもこのエラーがありました。コメントで指摘したにもかかわらず、修正するのを忘れていました。

さらに、この索引付けは私には正しく見えません。

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 に答える