1

我有以下计算两个数组的平方和误差的 CUDA 内核代码。

__global__ void kSquaredError(double* data, double* recon, double* error, 
                               unsigned int num_elements)
{
    const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    for (unsigned int i = idx; i < num_elements; i += blockDim.x * gridDim.x) {
        *error += pow(data[i] - recon[i], 2);
    }
}

我需要一个标量输出（错误）。在这种情况下，似乎所有线程都在同时写入错误。有什么方法需要同步吗？

目前我的结果很糟糕，所以我猜有一些问题。

4

1 回答 1

1

由于所有线程都尝试同时更新相同的全局内存地址，您现在正在执行的实现会受到竞争条件的影响。您可以轻松地放置一个atomicAdd函数，*error += pow...但由于每次更新都会对其进行序列化，因此会出现性能问题。

相反，您应该尝试使用共享内存进行减少，如下所示：

_global__ void kSquaredError(double* data, double* recon, double* error, unsigned int num_elements) {
    const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int tid = threadIdx.x;
    __shared__ double serror[blockDim.x];//temporary storage of each threads error

    for (unsigned int i = idx; i < num_elements; i += blockDim.x * gridDim.x) {
        serror[tid] += pow(data[i] - recon[i], 2);//put each threads value in shared memory
    }

    __syncthreads();

    int i = blockDim.x >> 1; //halve the threads
    for(;i>0;i>>=1) {//reduction in shared memory
            if(tid<i) {
                serror[tid] += serror[tid+i];
                __syncthreads();//make shure all threads are at the same state and have written to shared memory
            }
    }

    if(tid == 0) {//thread 0 updates the value in global memory
        atomicAdd(error,serror[tid]);// same as *error += serror[tid]; but atomic
    }
}

它的工作原理如下，每个线程都有自己的临时变量，它计算所有输入的错误总和，当它完成时，所有线程都聚集在__syncthreads指令处，以确保所有数据都是完整的。

现在，块中所有线程的一半将从对应的另一半中获取一个值，将其添加到自己的值中，再次一半的线程并再次执行，直到剩下一个线程（线程 0），该线程将具有总和.

现在线程 0 将使用 atomicAdd 函数更新全局内存，以避免与其他块（如果有的话）发生争用情况。

如果我们只使用第一个示例并在每个分配上使用 atomicAdd 。您将拥有gridDim.x*blockDim.x*num_elements可以序列化的原子函数，现在我们只有gridDim.x原子函数，数量要少得多。

请参阅优化 CUDA 中的并行缩减以进一步了解如何使用 cuda 进行缩减。

编辑

在减少 for 循环中添加了 if 语句，忘记了。

于 2013-01-21T20:56:58.933 回答