c++ - 迭代第二个循环，在 CUDA 中减少总和

Question

我必须将此代码从 C++ 并行化到 CUDA C

  for(ihist = 0; ihist < numhist; ihist++){ 
      for(iwin = 0; iwin<numwin; iwin++){
          denwham[ihist] += (numbinwin[iwin]/g[iwin])*exp(F[iwin]-U[ihist]); 
          }
          Punnorm[ihist] = numwham[ihist]/denwham[ihist];
        }

在 CUDA C 中，使用总和缩减:

extern __shared__ float sdata[];
  int tx = threadIdx.x;
  int i=blockIdx.x;
  int j=blockIdx.y;
  float sum=0.0;
  float temp=0.0;
  temp=U[j];


   if(tx<numwin)
   {
    sum=(numbinwin[tx]/g[tx])*exp(F[tx]- temp); 
    sdata[tx] = sum;
     __syncthreads();  
   }


  for(int offset = blockDim.x / 2;offset > 0;offset >>= 1)
  {
   if(tx < offset)
   {
    // add a partial sum upstream to our own
    sdata[tx] += sdata[tx + offset];
   }
   __syncthreads();
  }

   // finally, thread 0 writes the result
  if(threadIdx.x == 0)
  {
   // note that the result is per-block
   // not per-thread
   denwham[i] = sdata[0];

    for(int k=0;k<numhist;k++)
    Punnorm[k] = numwham[k]/denwham[k];
  }

并以这种方式初始化它:

 int smem_sz = (256)*sizeof(float);
  dim3 Block(numhist,numhist,1);
  NewProbabilitiesKernel<<<Block,256,smem_sz>>>(...);

我的问题是我无法使用 exp 遍历 U，我尝试了以下方法:

a) loop for/while inside the kernel that iterates over U 
b) iterate by thread
c) iterate to block

所有这些尝试导致我在 C++ 代码和代码 cuda 之间得到不同的结果。如果我输入一个常量而不是 U [i]，代码工作正常!

你有什么想法可以帮助我吗？

谢谢。

Answer 1

看起来你可以将 U 移出内循环

for(iwin = 0; iwin<numwin; iwin++){
    denwham += numbinwin[iwin] / g[iwin] * exp(F[iwin]); 
}
for(ihist = 0; ihist < numhist; ihist++){ 
    Punnorm[ihist] = numwham[ihist] / denwham * exp(U[ihist]);
}

更新

之后，您可以使用 2 个简单内核而不是 1 个复杂内核来完成任务。

1. 减少内核计算 denwham;
2. 用于计算Punnorm的一维变换核；