CUDA 并行化依赖的 2D 阵列

CUDA parallelizing a dependent 2D array

本文关键字：2D 阵列 依赖 并行化 CUDA      更新时间：2023-10-16

我有一个以下形式的示例循环。请注意，我的psi[i][j]依赖于psi[i+1][j], psi[i-1][j], psi[i][j+1] and psi[i][j-1]，我必须仅计算内部矩阵的psi。现在我尝试在 CUDA 中编写它，但结果与顺序不同。

for(i=1;i<=leni-2;i++)
for(j=1;j<=lenj-2;j++){
    psi[i][j]=(omega[i][j]*(dx*dx)*(dy*dy)+(psi[i+1][j]+psi[i-1][j])*(dy*dy)+(psi[i][j+1]+psi[i][j-1])*(dx*dx) )/(2.0*(dx*dx)+2.0*(dy*dy));
}

这是我的 CUDA 格式。

//KERNEL
__global__ void ComputePsi(double *psi, double *omega, int imax, int jmax)
{
int x = blockIdx.x;
int y = blockIdx.y;
int i = (jmax*x) + y;
double beta = 1;
double dx=(double)30/(imax-1);
double dy=(double)1/(jmax-1);
if((i)%jmax!=0 && (i+1)%jmax!=0 &&  i>=jmax && i<imax*jmax-jmax){
    psi[i]=(omega[i]*(dx*dx)*(dy*dy)+(psi[i+jmax]+psi[i-jmax])*(dy*dy)+(psi[i+1]+psi[i-1])*(dx*dx) )/(2.0*(dx*dx)+2.0*(dy*dy));
}
}

//Code
cudaMalloc((void **) &dev_psi, leni*lenj*sizeof(double));
cudaMalloc((void **) &dev_omega, leni*lenj*sizeof(double));
cudaMemcpy(dev_psi, psi, leni*lenj*sizeof(double),cudaMemcpyHostToDevice);
cudaMemcpy(dev_omega, omega, leni*lenj*sizeof(double),cudaMemcpyHostToDevice);
dim3 grids(leni,lenj);
for(iterpsi=0;iterpsi<30;iterpsi++)         
    ComputePsi<<<grids,1>>>(dev_psi, dev_omega, leni, lenj);

其中psi[leni][lenj] and omega[leni][lenj]和双精度数组。

问题是顺序的，CUDA 代码给出不同的结果。代码中是否需要任何修改？

您正在使用全局内存，并且正在更改 psi 条目，而其他线程可能需要旧值。只需将新迭代的值存储在单独的变量中即可。但请记住，您必须在每次迭代后交换变量！更复杂的方法是将共享内存和空间域分配给单独的线程的解决方案。只需谷歌搜索用于求解热/扩散方程的 CUDA 教程，您就会明白。

for(i=1;i<=leni-2;i++)
  for(j=1;j<=lenj-2;j++){
    psi[i][j]= ( omega[i][j]*(dx*dx)*(dy*dy) + 
                 (psi[i+1][j]+psi[i-1][j]) * (dy*dy) + 
                 (psi[i][j+1]+psi[i][j-1]) * (dx*dx) 
               )/(2.0*(dx*dx)+2.0*(dy*dy));
}

我认为这个内核的顺序也不正确：psi[i][j]的值取决于此处操作的顺序 - 因此您将使用未更新的psi[i+1][j]和psi[i][j+1]，但psi[i-1][j]和psi[i][j-1]已在此扫描中更新。

确保使用 CUDA 的结果会有所不同，其中操作的顺序不同。

要强制执行这样的排序，如果可能的话，您需要插入如此多的同步，这对于 CUDA 来说可能不值得。这真的是你需要做的吗？

相关文章：