将“for”循环转换为CUDA并行代码

我正在转换CUDA并行化代码中最初编写的C代码的过程。还是一个新手，我把大部分代码转换成了CUDA，但是我的一些内核没有正确地完成这项工作。将“for”循环转换为CUDA并行代码

这里是我的内核：

__global__ void kernel(long int *neighbour, double *f, double *r, double *b, double *fn, double *rn, double *bn, int nfluidsite){ 

int ns = blockDim.x * blockIdx.x + threadIdx.x; 

    if(ns<nfluidsite) 
{ 
    double tempr = r[ns]; 
    double tempb = b[ns]; 
    rn[ns]=tempr; 
    bn[ns]=tempb; 
    for(int q=1;q<Q;++q) 
    { 
    double confr=r[q*NSITE+ns]; 
    double confb=b[q*NSITE+ns]; 
    __syncthreads(); 
    int ns1=neighbour[q*NTOTAL+ns]; 
    __syncthreads(); 
    rn[q*NSITE+ns1]=confr; 
    bn[q*NSITE+ns1]=confb; 
    } 
} 

if(ns<NSITE) 
{ 
    for(int q=0;q<Q;++q) 
    { 
     double rqns = rn[q*NSITE+ns]; 
     double bqns = bn[q*NSITE+ns]; 
    __syncthreads(); 
    r[q*NSITE+ns]=rqns; 
    b[q*NSITE+ns]=bqns; 
    f[q*NSITE+ns]=rqns+bqns; 
    } 
} 

}

所以，这个代码工作正常（虽然它不是在所有优化），但我也希望进行并行内的循环上q。所以，我是这样的：

int ns = blockIdx.x; 
    int q = threadIdx.x;

，我开始了我的内核如下：

blocksPerGrid = NSITE; 
threadsPerBlock = Q; 
kernel<<<blocksPerGrid,threadsPerBlock>>>(neighbourCu, fCu, rCu, bCu, fnCu, rnCu, bnCu, nfluidsite);

而且它不会在所有的工作，CUDA不会对数组返回任何错误，但操作是随机的...我在完全并行版本中添加了__syncthreads()命令，但它并没有解决这些差异。

而且，我不为什么，但如果我使用超过1024个线程，在我的内核中的指令也运行随机...

嗯，我一直不解两周内，如果有人看到我需要做的，请给我一个提示！

来源

2014-02-27 Seif

我认为你不需要__syncthreads();你已经放置它们，你正在分配和读取线程本地（寄存器或本地内存）变量。

在另一边肯定需要调用__threadfence();的代码块前开始：

if(ns<NSITE) { 
    ... 
}

的违规声明：

rn[q*NSITE+ns1]=confr; 
    bn[q*NSITE+ns1]=confb;

前：

double rqns = rn[q*NSITE+ns]; 
    double bqns = bn[q*NSITE+ns];

之后。

由于一个线程写入一些全局内存和一个不同读取它，你需要至少两个放在中间：

__threadfence(); 
__syncthreads();

（更多信息here）。只有rn和bn被同一个块中的线程修改，这才能正确工作。如果碰巧块外部的线程可以修改它们，这还不够：在继续之前，需要保证所有块都已经到达该点（注意：__syncthreads()只保证同一块内的线程已经达到了这一点 - 它是块局部障碍）。你有三种选择：

在这一点上将内核拆分到两个不同的内核中 - 单独的内核调用被隐式地同步。我会从这开始，看看它是否有效。

如果NSITE只是一个线程块，就像看起来一样，您可以再次遵循全局总和here的示例：只有最后一个块会执行 - 其余所有其他块都会跳过。

在这一点上实现了一个全球性障碍（这不是一件容易的事情 - 但如果你在谷歌周围，你会发现很好的参考）。

来源

2014-02-27 11:27:09 Sigismondo

谢谢！我使用了第一个选项（最简单的一个），它工作。它也解决了我的线程每块数限制问题。如果在其他地方需要，我会记住其他选项。 – Seif

将“for”循环转换为CUDA并行代码

回答

相关问题