0

我在 Linux 系统和 Tesla C2075 机器上运行。我正在启动一个内核，它是缩减内核的修改版本。我的目标是找到大型数据集（结果）的平均值和逐步平均版本（time_avg）。请参阅下面的代码。

“result”和“time_avg”的大小相同，等于“nsamps”。“time_avg”包含数组结果的连续平均组。因此，前半部分包含每两个非重叠样本的平均值，之后的四分之一包含每四个非重叠样本的平均值，接下来的八分之一样本包含 8 个样本，依此类推。

__global__ void timeavg_mean(float *result, unsigned int *nsamps, float *time_avg, float *mean) {

__shared__ float temp[1024];
int ltid = threadIdx.x, gtid = blockIdx.x*blockDim.x + threadIdx.x, stride;
int start = 0, index;
unsigned int npts = *nsamps;

printf("here here\n");

// Store chunk of memory=2*blockDim.x (which is to be reduced) into shared memory
if ( (2*gtid) < npts ){
        temp[2*ltid] = result[2*gtid];
        temp[2*ltid+1] = result[2*gtid + 1]; 
}    
        for (stride=1; stride<blockDim.x; stride>>=1) {
                __syncthreads();
                if (ltid % (stride*2) == 0){ 
                  if ( (2*gtid) < npts ){
                        temp[2*ltid] += temp[2*ltid + stride];
                        index = (int)(start + gtid/stride);
                        time_avg[index] = (float)( temp[2*ltid]/(2.0*stride) );
                  }   
                }   
        start += npts/(2*stride);
        }   
__syncthreads();
if (ltid == 0)
{
        atomicAdd(mean, temp[0]);    
}
__syncthreads();
printf("%f\n", *mean);
}

启动配置为 40 个块，512 个线程。数据集约为 40k 个样本。

在我的主代码中，我cudaGetLastError()在内核调用之后调用它并且它没有返回错误。内存分配和内存副本不返回错误。如果我在内核调用之后写入cudaDeviceSynchronize()（或 acudaMemcpy检查均值的值），则程序在内核调用之后完全挂起。如果我删除它，程序运行并退出。在这两种情况下，我都得到“这里”的输出还是打印的平均值。我知道除非内核成功执行，否则 printf 不会打印。这与__syncthreads()递归有关吗？所有线程都会达到相同的深度，所以我认为可以检查。

这里有什么问题？

谢谢！

4

2 回答 2

4

内核调用是异步的，如果内核成功启动，您的主机代码将继续运行并且您不会看到任何错误。内核运行期间发生的错误仅在您执行显式同步或调用导致隐式同步的函数后才会出现。
如果您的主机在同步时挂起，那么您的内核可能还没有完成运行——它要么正在运行某个无限循环，要么正在等待某个__synchthreads()或其他同步原语。
您的代码似乎包含一个无限循环：for (stride=1; stride<blockDim.x; stride>>=1). 您可能想要向左而不是向右移动步幅：stride<<=1.

您提到了递归，但您的代码仅包含一个__global__函数，没有递归调用。

于 2013-03-24T13:39:42.377 回答
1

你的内核有一个无限循环。将 for 循环替换为

for (stride=1; stride<blockDim.x; stride<<=1) {
于 2013-03-24T13:51:41.707 回答