CUDA C++共享内存和 if 条件

Question

我有一个问题，我找不到自己的答案，我希望你们中的一些人能为我提供一些关于可能的解决方案的见解。在内核调用中，我想插入一个关于访问共享内存的 if 条件。

__global__ void GridFillGPU (int * gridGLOB, int n) {
    __shared__ int grid[SIZE] // ... initialized to zero
    int tid = threadIdx.x
        if (tid < n) {
            for ( int k = 0; k < SIZE; k++) { 
                if (grid[k] == 0) {
                    grid[k] = tid+1;
                    break;
                }
            }
        }
    //... here write grid to global memory gridGLOB
    }

这个想法是，如果元素 grid[k] 已经由一个线程（带有索引 tid）写入，则不应由另一个线程写入。我的问题是：这甚至可以并行完成吗？由于所有并行线程都执行相同的 for 循环，我如何确保正确计算 if 条件？我猜这会导致某些比赛条件。我对 Cuda 很陌生，所以我希望这个问题不愚蠢。我知道网格需要在共享内存中，并且应该避免if语句，但目前我没有找到其他方法。我很感激任何帮助

编辑：这是显式版本，它解释了为什么数组被称为网格

__global__ void GridFillGPU (int * pos, int * gridGLOB, int n) {
    __shared__ int grid[SIZE*7] // ... initialized to zero
    int tid = threadIdx.x
        if (tid < n) {
        int jmin = pos[tid] - 3;
        int jmax = pos[tid] + 3;
          for ( int j = jmin; j <= jmax; j++ { 
            for ( int k = 0; k < SIZE; k++) { 
                if (grid[(j-jmin)*SIZE + k] == 0) {
                    grid[(j-jmin)*SIZE + k] = tid+1;
                    break;
                }
            }
        }
    } //... here write grid to global memory gridGLOB
}

Answer 1

你应该以一种你不需要担心"如果已经写好"的方式对你的问题进行建模，也因为 cuda 不保证线程的执行顺序，所以顺序可能不是你执行的方式。有一些小事情可以确保您在经线中明智地订购，但事实并非如此。您可以使用同步栏和东西，但我认为不是您的情况。

如果你正在处理一个网格，你应该以一种方式建模，即每个线程都有自己的内存区域。 并且不应该与其他线程区域重叠（至少在写作中，在阅读中你可以超越边界）。此外，我不会担心共享内存，首先使算法工作，然后考虑优化，例如使用扭曲在共享内存中加载图块。

在这种情况下，如果你想在网格中拆分你的域，你应该设置内核，以便有足够的线程作为你的网格"单元格"或像素（如果是图像）。然后，您使用 cuda 为您提供的线程和块坐标来计算您应该在内存中读取和写入的位置。

关于 cuda 的 udacity.com 有一个非常好的课程，你可能想看看。https://www.udacity.com/courses/cs344coursera.com 还有另一个，但我不知道它现在是否开放。无论如何，在网格中划分域是一个非常常见且已解决的问题，您可以找到很多材料。