Cuda 内核中的数据组织

Question

我是 Cuda 的新手，一直在阅读教程和其他开源代码，以尝试理解事物。我知道线程层次结构的一般概念。

TL;DR，我阅读的所有教程都假设发送到内核的数据也按此层次结构组织，而在启动内核之前没有明确这样做。传递给内核的数据不应该在传递给内核之前重新排列在网格>块>线程层次结构中吗？以下是在这方面让我感到困惑的两个 sn-ps。

我遵循了这个x_plus_y 教程here。在本教程中，sn-p 如下：

_global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

在上面的sn-p中，我们想在x和y中添加对应的元素，但是我怎么知道，x和y在GPU中是这样放置的，索引@ 987654328@（使用blockIdx、blockDim等计算）实际上指向x和y的对应元素。如果x 和y 在内存中一个接一个地放置，那么用于y 的索引不应该考虑x 的长度吗？我在这里缺少一些关键的直观理解。另外，我怎么知道在 GPU 中的哪里映射了数组的一些随机元素，比如 x[1011]？还是由于某种抽象，我不需要关心数据的显式定位？

我还将讨论另一个 sn-p，来自一个开源的 torch repo。这是用于计算两组点云之间距离度量的内核。每个云都是一个Nx3 矩阵（具有N 3-D 点）。

b 是批量大小（因此，b 的云数量传递给内核）

n是第一组每个云中的点数

m 是第二组每个云中的点数。

例如，第一组云可以是 (16,1024,3) 和第二组 (16,512,3)：

__global__ void NmDistanceKernel(int b,int n,const float * xyz,int m,const float * xyz2,float * result,int * result_i){
    const int batch=512;
    __shared__ float buf[batch*3];
    for (int i=blockIdx.x;i<b;i+=gridDim.x){
        for (int k2=0;k2<m;k2+=batch){
            int end_k=min(m,k2+batch)-k2;
            for (int j=threadIdx.x;j<end_k*3;j+=blockDim.x){
                buf[j]=xyz2[(i*m+k2)*3+j];
            }


        for (int j=threadIdx.x+blockIdx.y*blockDim.x;j<n;j+=blockDim.x*gridDim.y){
                float x1=xyz[(i*n+j)*3+0];
                float y1=xyz[(i*n+j)*3+1];
                float z1=xyz[(i*n+j)*3+2];
            }
    }
}

以上内核，启动如下：

NmDistanceKernel<<<dim3(32,16,1),512>>>(batch_size, n, xyz1.data<float>(), m, xyz2.data<float>(), dist1.data<float>(), idx1.data<int>());

同样，在上面的内核中，作者假设他们传递给内核的数据是经过组织的，以便索引机制可以工作。他们没有明确地将每个点放在每个线程中，然后将一堆点放在一个块内，将一堆云放在一个网格内。然而，这种结构是在内核内部假设的。

Answer 1

在调用内核之前，您必须已将数据放入 GPU。

数据大多以数据数组的形式传入，因此这些数组的结构在 GPU 上与在主机代码中的结构相同。

在第一个示例中，数组 x 和 y 分别传入，因此 x 和 y 的索引都从 0 开始。您可以将它们传递到一个大数组中，然后需要调整索引。

这已在您的其他示例中完成。数组xyz 由所有点的 x y 和 z 值组成。订单类似于x1 y1 z1 x2 y2 z2 x3 y3 z3 ...。这就是为什么在访问您看到x = [...]+0; y = [...]+1; z = [...]+2; 的值时。对于下一点，指数都增加了 3。

要访问内核中的数据，您需要参考 CUDA 提供的标识符。您使用网格和块中线程的位置。

在第一个示例中，程序员选择启动所有读取数组中第一个连续条目的线程。他通过为每个线程分配一个唯一的index 来做到这一点：

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

threadIdx.x 告诉我们线程驻留在块中的位置，所以如果我们只启动一个块就足够了。但是不同块中的不同线程将具有相同的索引。我们必须通过获取他们的blockIdx.x 来将它们分开。该块的长度为blockDim.x，第二个块中的第一个线程应该在块1中的最后一个线程之后继续。所以上面的index公式形成。

然后，每个线程向前跳转，以便第一个线程接下来读取最后一个线程刚刚读取的数据之后的第一个数据，依此类推。

您启动的网格使用的维度越多，这些计算就必须越复杂。如果您对它们感到满意，请尝试从简单的网格开始并增加复杂性。